[mlir] take MLIRContext instead of LLVMDialect in getters of LLVMType's

[ftynse]

Summary

Historical modeling of the LLVM dialect types had been wrapping LLVM IR types
and therefore needed access to the instance of LLVMContext stored in the
LLVMDialect. The new modeling does not rely on that and only needs the
MLIRContext that is used for uniquing, similarly to other MLIR types. Change
LLVMType::get<Kind>Ty functions to take MLIRContext * instead of
LLVMDialect * as first argument. This brings the code base closer to
completely removing the dependence on LLVMContext from the LLVMDialect,
together with additional support for thread-safety of its use.

Depends On D85371

mlir/examples/toy/Ch6/mlir/LowerToLLVM.cpp

explicit PrintOpLowering(MLIRContext *context)
       : ConversionPattern(toy::PrintOp::getOperationName(), 1, context) {}
 
   LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto memRefType = (*op->operand_type_begin()).cast<MemRefType>();
     auto memRefShape = memRefType.getShape();
     auto loc = op->getLoc();
-    auto *llvmDialect =
-        op->getContext()->getRegisteredDialect<LLVM::LLVMDialect>();
-    assert(llvmDialect && "expected llvm dialect to be registered");
 
     ModuleOp parentModule = op->getParentOfType<ModuleOp>();
 
     // Get a symbol reference to the printf function, inserting it if necessary.
-    auto printfRef = getOrInsertPrintf(rewriter, parentModule, llvmDialect);
+    auto printfRef = getOrInsertPrintf(rewriter, parentModule);
     Value formatSpecifierCst = getOrCreateGlobalString(
-        loc, rewriter, "frmt_spec", StringRef("%f \0", 4), parentModule,
-        llvmDialect);
+        loc, rewriter, "frmt_spec", StringRef("%f \0", 4), parentModule);
     Value newLineCst = getOrCreateGlobalString(
-        loc, rewriter, "nl", StringRef("\n\0", 2), parentModule, llvmDialect);
+        loc, rewriter, "nl", StringRef("\n\0", 2), parentModule);
 
     // Create a loop for each of the dimensions within the shape.
     SmallVector<Value, 4> loopIvs;
     for (unsigned i = 0, e = memRefShape.size(); i != e; ++i) {
       auto lowerBound = rewriter.create<ConstantIndexOp>(loc, 0);
       auto upperBound = rewriter.create<ConstantIndexOp>(loc, memRefShape[i]);
       auto step = rewriter.create<ConstantIndexOp>(loc, 1);
       auto loop =
matchAndRewrite(Operation *op, ArrayRef<Value> operands,
     rewriter.eraseOp(op);
     return success();
   }
 
 private:
   /// Return a symbol reference to the printf function, inserting it into the
   /// module if necessary.
   static FlatSymbolRefAttr getOrInsertPrintf(PatternRewriter &rewriter,
-                                             ModuleOp module,
-                                             LLVM::LLVMDialect *llvmDialect) {
+                                             ModuleOp module) {
     auto *context = module.getContext();
     if (module.lookupSymbol<LLVM::LLVMFuncOp>("printf"))
       return SymbolRefAttr::get("printf", context);
 
     // Create a function declaration for printf, the signature is:
     //   * `i32 (i8*, ...)`
-    auto llvmI32Ty = LLVM::LLVMType::getInt32Ty(llvmDialect);
-    auto llvmI8PtrTy = LLVM::LLVMType::getInt8PtrTy(llvmDialect);
+    auto llvmI32Ty = LLVM::LLVMType::getInt32Ty(context);
+    auto llvmI8PtrTy = LLVM::LLVMType::getInt8PtrTy(context);
     auto llvmFnType = LLVM::LLVMType::getFunctionTy(llvmI32Ty, llvmI8PtrTy,
                                                     /*isVarArg=*/true);
 
     // Insert the printf function into the body of the parent module.
     PatternRewriter::InsertionGuard insertGuard(rewriter);
     rewriter.setInsertionPointToStart(module.getBody());
     rewriter.create<LLVM::LLVMFuncOp>(module.getLoc(), "printf", llvmFnType);
     return SymbolRefAttr::get("printf", context);
   }
 
   /// Return a value representing an access into a global string with the given
   /// name, creating the string if necessary.
   static Value getOrCreateGlobalString(Location loc, OpBuilder &builder,
                                        StringRef name, StringRef value,
-                                       ModuleOp module,
-                                       LLVM::LLVMDialect *llvmDialect) {
+                                       ModuleOp module) {
     // Create the global at the entry of the module.
     LLVM::GlobalOp global;
     if (!(global = module.lookupSymbol<LLVM::GlobalOp>(name))) {
       OpBuilder::InsertionGuard insertGuard(builder);
       builder.setInsertionPointToStart(module.getBody());
       auto type = LLVM::LLVMType::getArrayTy(
-          LLVM::LLVMType::getInt8Ty(llvmDialect), value.size());
+          LLVM::LLVMType::getInt8Ty(builder.getContext()), value.size());
       global = builder.create<LLVM::GlobalOp>(loc, type, /*isConstant=*/true,
                                               LLVM::Linkage::Internal, name,
                                               builder.getStringAttr(value));
     }
 
     // Get the pointer to the first character in the global string.
     Value globalPtr = builder.create<LLVM::AddressOfOp>(loc, global);
     Value cst0 = builder.create<LLVM::ConstantOp>(
-        loc, LLVM::LLVMType::getInt64Ty(llvmDialect),
+        loc, LLVM::LLVMType::getInt64Ty(builder.getContext()),
         builder.getIntegerAttr(builder.getIndexType(), 0));
     return builder.create<LLVM::GEPOp>(
-        loc, LLVM::LLVMType::getInt8PtrTy(llvmDialect), globalPtr,
+        loc, LLVM::LLVMType::getInt8PtrTy(builder.getContext()), globalPtr,
         ArrayRef<Value>({cst0, cst0}));
   }
 };
 } // end anonymous namespace
 
 //===----------------------------------------------------------------------===//
 // ToyToLLVMLoweringPass
 //===----------------------------------------------------------------------===//

mlir/examples/toy/Ch7/mlir/LowerToLLVM.cpp

explicit PrintOpLowering(MLIRContext *context)
       : ConversionPattern(toy::PrintOp::getOperationName(), 1, context) {}
 
   LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto memRefType = (*op->operand_type_begin()).cast<MemRefType>();
     auto memRefShape = memRefType.getShape();
     auto loc = op->getLoc();
-    auto *llvmDialect =
-        op->getContext()->getRegisteredDialect<LLVM::LLVMDialect>();
-    assert(llvmDialect && "expected llvm dialect to be registered");
 
     ModuleOp parentModule = op->getParentOfType<ModuleOp>();
 
     // Get a symbol reference to the printf function, inserting it if necessary.
-    auto printfRef = getOrInsertPrintf(rewriter, parentModule, llvmDialect);
+    auto printfRef = getOrInsertPrintf(rewriter, parentModule);
     Value formatSpecifierCst = getOrCreateGlobalString(
-        loc, rewriter, "frmt_spec", StringRef("%f \0", 4), parentModule,
-        llvmDialect);
+        loc, rewriter, "frmt_spec", StringRef("%f \0", 4), parentModule);
     Value newLineCst = getOrCreateGlobalString(
-        loc, rewriter, "nl", StringRef("\n\0", 2), parentModule, llvmDialect);
+        loc, rewriter, "nl", StringRef("\n\0", 2), parentModule);
 
     // Create a loop for each of the dimensions within the shape.
     SmallVector<Value, 4> loopIvs;
     for (unsigned i = 0, e = memRefShape.size(); i != e; ++i) {
       auto lowerBound = rewriter.create<ConstantIndexOp>(loc, 0);
       auto upperBound = rewriter.create<ConstantIndexOp>(loc, memRefShape[i]);
       auto step = rewriter.create<ConstantIndexOp>(loc, 1);
       auto loop =
matchAndRewrite(Operation *op, ArrayRef<Value> operands,
     rewriter.eraseOp(op);
     return success();
   }
 
 private:
   /// Return a symbol reference to the printf function, inserting it into the
   /// module if necessary.
   static FlatSymbolRefAttr getOrInsertPrintf(PatternRewriter &rewriter,
-                                             ModuleOp module,
-                                             LLVM::LLVMDialect *llvmDialect) {
+                                             ModuleOp module) {
     auto *context = module.getContext();
     if (module.lookupSymbol<LLVM::LLVMFuncOp>("printf"))
       return SymbolRefAttr::get("printf", context);
 
     // Create a function declaration for printf, the signature is:
     //   * `i32 (i8*, ...)`
-    auto llvmI32Ty = LLVM::LLVMType::getInt32Ty(llvmDialect);
-    auto llvmI8PtrTy = LLVM::LLVMType::getInt8PtrTy(llvmDialect);
+    auto llvmI32Ty = LLVM::LLVMType::getInt32Ty(context);
+    auto llvmI8PtrTy = LLVM::LLVMType::getInt8PtrTy(context);
     auto llvmFnType = LLVM::LLVMType::getFunctionTy(llvmI32Ty, llvmI8PtrTy,
                                                     /*isVarArg=*/true);
 
     // Insert the printf function into the body of the parent module.
     PatternRewriter::InsertionGuard insertGuard(rewriter);
     rewriter.setInsertionPointToStart(module.getBody());
     rewriter.create<LLVM::LLVMFuncOp>(module.getLoc(), "printf", llvmFnType);
     return SymbolRefAttr::get("printf", context);
   }
 
   /// Return a value representing an access into a global string with the given
   /// name, creating the string if necessary.
   static Value getOrCreateGlobalString(Location loc, OpBuilder &builder,
                                        StringRef name, StringRef value,
-                                       ModuleOp module,
-                                       LLVM::LLVMDialect *llvmDialect) {
+                                       ModuleOp module) {
     // Create the global at the entry of the module.
     LLVM::GlobalOp global;
     if (!(global = module.lookupSymbol<LLVM::GlobalOp>(name))) {
       OpBuilder::InsertionGuard insertGuard(builder);
       builder.setInsertionPointToStart(module.getBody());
       auto type = LLVM::LLVMType::getArrayTy(
-          LLVM::LLVMType::getInt8Ty(llvmDialect), value.size());
+          LLVM::LLVMType::getInt8Ty(builder.getContext()), value.size());
       global = builder.create<LLVM::GlobalOp>(loc, type, /*isConstant=*/true,
                                               LLVM::Linkage::Internal, name,
                                               builder.getStringAttr(value));
     }
 
     // Get the pointer to the first character in the global string.
     Value globalPtr = builder.create<LLVM::AddressOfOp>(loc, global);
     Value cst0 = builder.create<LLVM::ConstantOp>(
-        loc, LLVM::LLVMType::getInt64Ty(llvmDialect),
+        loc, LLVM::LLVMType::getInt64Ty(builder.getContext()),
         builder.getIntegerAttr(builder.getIndexType(), 0));
     return builder.create<LLVM::GEPOp>(
-        loc, LLVM::LLVMType::getInt8PtrTy(llvmDialect), globalPtr,
+        loc, LLVM::LLVMType::getInt8PtrTy(builder.getContext()), globalPtr,
         ArrayRef<Value>({cst0, cst0}));
   }
 };
 } // end anonymous namespace
 
 //===----------------------------------------------------------------------===//
 // ToyToLLVMLoweringPass
 //===----------------------------------------------------------------------===//

mlir/include/mlir/Dialect/LLVMIR/LLVMDialect.h

 
 #include "mlir/Dialect/LLVMIR/LLVMOpsDialect.h.inc"
 
 /// Create an LLVM global containing the string "value" at the module containing
 /// surrounding the insertion point of builder. Obtain the address of that
 /// global and use it to compute the address of the first character in the
 /// string (operations inserted at the builder insertion point).
 Value createGlobalString(Location loc, OpBuilder &builder, StringRef name,
-                         StringRef value, LLVM::Linkage linkage,
-                         LLVM::LLVMDialect *llvmDialect);
+                         StringRef value, LLVM::Linkage linkage);
 
 /// LLVM requires some operations to be inside of a Module operation. This
 /// function confirms that the Operation has the desired properties.
 bool satisfiesLLVMModule(Operation *op);
 
 /// Clones the given module into the provided context. This is implemented by
 /// transforming the module into bitcode and then reparsing the bitcode in the
 /// provided context.
 std::unique_ptr<llvm::Module>
 cloneModuleIntoNewContext(llvm::LLVMContext *context, llvm::Module *module);
 
 } // end namespace LLVM
 } // end namespace mlir
 
 #endif // MLIR_DIALECT_LLVMIR_LLVMDIALECT_H_

mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td

 // Integer type of a specific width.
 class LLVMI<int width>
     : Type<And<[
         LLVM_Type.predicate,
         CPred<
          "$_self.cast<::mlir::LLVM::LLVMType>().isIntegerTy(" # width # ")">]>,
          "LLVM dialect " # width # "-bit integer">,
       BuildableType<
-        "::mlir::LLVM::LLVMType::getIntNTy("
-        "$_builder.getContext()->getRegisteredDialect<LLVM::LLVMDialect>(),"
+        "::mlir::LLVM::LLVMType::getIntNTy($_builder.getContext(),"
          # width # ")">;
 
 def LLVMI1 : LLVMI<1>;
 
 // Base class for LLVM operations. Defines the interface to the llvm::IRBuilder
 // used to translate to LLVM IR proper.
 class LLVM_OpBase<Dialect dialect, string mnemonic, list<OpTrait> traits = []> :
     Op<dialect, mnemonic, traits> {

mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td

def LLVM_ICmpOp : LLVM_OneResultOp<"icmp", [NoSideEffect]>,
                   Arguments<(ins ICmpPredicate:$predicate, LLVM_Type:$lhs,
                              LLVM_Type:$rhs)> {
   let llvmBuilder = [{
     $res = builder.CreateICmp(getLLVMCmpPredicate($predicate), $lhs, $rhs);
   }];
   let builders = [OpBuilder<
     "OpBuilder &b, OperationState &result, ICmpPredicate predicate, Value lhs, "
     "Value rhs", [{
-      LLVMDialect *dialect = &lhs.getType().cast<LLVMType>().getDialect();
-      build(b, result, LLVMType::getInt1Ty(dialect),
+      build(b, result, LLVMType::getInt1Ty(lhs.getType().getContext()),
             b.getI64IntegerAttr(static_cast<int64_t>(predicate)), lhs, rhs);
     }]>];
   let parser = [{ return parseCmpOp<ICmpPredicate>(parser, result); }];
   let printer = [{ printICmpOp(p, *this); }];
 }
 
 // Predicate for float comparisons
 def FCmpPredicateFALSE  : I64EnumAttrCase<"_false", 0>;
def LLVM_FCmpOp : LLVM_OneResultOp<"fcmp", [NoSideEffect]>,
                   Arguments<(ins FCmpPredicate:$predicate, LLVM_Type:$lhs,
                              LLVM_Type:$rhs)> {
   let llvmBuilder = [{
     $res = builder.CreateFCmp(getLLVMCmpPredicate($predicate), $lhs, $rhs);
   }];
   let builders = [OpBuilder<
     "OpBuilder &b, OperationState &result, FCmpPredicate predicate, Value lhs, "
     "Value rhs", [{
-      LLVMDialect *dialect = &lhs.getType().cast<LLVMType>().getDialect();
-      build(b, result, LLVMType::getInt1Ty(dialect),
+      build(b, result, LLVMType::getInt1Ty(lhs.getType().getContext()),
             b.getI64IntegerAttr(static_cast<int64_t>(predicate)), lhs, rhs);
     }]>];
   let parser = [{ return parseCmpOp<FCmpPredicate>(parser, result); }];
   let printer = [{ printFCmpOp(p, *this); }];
 }
 
 // Floating point binary operations.
 def LLVM_FAddOp : LLVM_ArithmeticOp<"fadd", "CreateFAdd">;

mlir/include/mlir/Dialect/LLVMIR/LLVMTypes.h

public:
   static bool isValidPointerElementType(LLVMType type);
 
   /// Struct type utilities.
   LLVMType getStructElementType(unsigned i);
   unsigned getStructNumElements();
   bool isStructTy();
 
   /// Utilities used to generate floating point types.
-  static LLVMType getDoubleTy(LLVMDialect *dialect);
-  static LLVMType getFloatTy(LLVMDialect *dialect);
-  static LLVMType getBFloatTy(LLVMDialect *dialect);
-  static LLVMType getHalfTy(LLVMDialect *dialect);
-  static LLVMType getFP128Ty(LLVMDialect *dialect);
-  static LLVMType getX86_FP80Ty(LLVMDialect *dialect);
+  static LLVMType getDoubleTy(MLIRContext *context);
+  static LLVMType getFloatTy(MLIRContext *context);
+  static LLVMType getBFloatTy(MLIRContext *context);
+  static LLVMType getHalfTy(MLIRContext *context);
+  static LLVMType getFP128Ty(MLIRContext *context);
+  static LLVMType getX86_FP80Ty(MLIRContext *context);
 
   /// Utilities used to generate integer types.
-  static LLVMType getIntNTy(LLVMDialect *dialect, unsigned numBits);
-  static LLVMType getInt1Ty(LLVMDialect *dialect) {
-    return getIntNTy(dialect, /*numBits=*/1);
+  static LLVMType getIntNTy(MLIRContext *context, unsigned numBits);
+  static LLVMType getInt1Ty(MLIRContext *context) {
+    return getIntNTy(context, /*numBits=*/1);
   }
-  static LLVMType getInt8Ty(LLVMDialect *dialect) {
-    return getIntNTy(dialect, /*numBits=*/8);
+  static LLVMType getInt8Ty(MLIRContext *context) {
+    return getIntNTy(context, /*numBits=*/8);
   }
-  static LLVMType getInt8PtrTy(LLVMDialect *dialect) {
-    return getInt8Ty(dialect).getPointerTo();
+  static LLVMType getInt8PtrTy(MLIRContext *context) {
+    return getInt8Ty(context).getPointerTo();
   }
-  static LLVMType getInt16Ty(LLVMDialect *dialect) {
-    return getIntNTy(dialect, /*numBits=*/16);
+  static LLVMType getInt16Ty(MLIRContext *context) {
+    return getIntNTy(context, /*numBits=*/16);
   }
-  static LLVMType getInt32Ty(LLVMDialect *dialect) {
-    return getIntNTy(dialect, /*numBits=*/32);
+  static LLVMType getInt32Ty(MLIRContext *context) {
+    return getIntNTy(context, /*numBits=*/32);
   }
-  static LLVMType getInt64Ty(LLVMDialect *dialect) {
-    return getIntNTy(dialect, /*numBits=*/64);
+  static LLVMType getInt64Ty(MLIRContext *context) {
+    return getIntNTy(context, /*numBits=*/64);
   }
 
   /// Utilities used to generate other miscellaneous types.
   static LLVMType getArrayTy(LLVMType elementType, uint64_t numElements);
   static LLVMType getFunctionTy(LLVMType result, ArrayRef<LLVMType> params,
                                 bool isVarArg);
   static LLVMType getFunctionTy(LLVMType result, bool isVarArg) {
     return getFunctionTy(result, llvm::None, isVarArg);
   }
-  static LLVMType getStructTy(LLVMDialect *dialect, ArrayRef<LLVMType> elements,
+  static LLVMType getStructTy(MLIRContext *context, ArrayRef<LLVMType> elements,
                               bool isPacked = false);
-  static LLVMType getStructTy(LLVMDialect *dialect, bool isPacked = false) {
-    return getStructTy(dialect, llvm::None, isPacked);
+  static LLVMType getStructTy(MLIRContext *context, bool isPacked = false) {
+    return getStructTy(context, llvm::None, isPacked);
   }
   template <typename... Args>
   static typename std::enable_if<llvm::are_base_of<LLVMType, Args...>::value,
                                  LLVMType>::type
   getStructTy(LLVMType elt1, Args... elts) {
     SmallVector<LLVMType, 8> fields({elt1, elts...});
-    return getStructTy(&elt1.getDialect(), fields);
+    return getStructTy(elt1.getContext(), fields);
   }
   static LLVMType getVectorTy(LLVMType elementType, unsigned numElements);
 
   /// Void type utilities.
-  static LLVMType getVoidTy(LLVMDialect *dialect);
+  static LLVMType getVoidTy(MLIRContext *context);
   bool isVoidTy();
 
   // Creation and setting of LLVM's identified struct types
-  static LLVMType createStructTy(LLVMDialect *dialect,
+  static LLVMType createStructTy(MLIRContext *context,
                                  ArrayRef<LLVMType> elements,
                                  Optional<StringRef> name,
                                  bool isPacked = false);
 
-  static LLVMType createStructTy(LLVMDialect *dialect,
+  static LLVMType createStructTy(MLIRContext *context,
                                  Optional<StringRef> name) {
-    return createStructTy(dialect, llvm::None, name);
+    return createStructTy(context, llvm::None, name);
   }
 
   static LLVMType createStructTy(ArrayRef<LLVMType> elements,
                                  Optional<StringRef> name,
                                  bool isPacked = false) {
     assert(!elements.empty() &&
            "This method may not be invoked with an empty list");
     LLVMType ele0 = elements.front();
-    return createStructTy(&ele0.getDialect(), elements, name, isPacked);
+    return createStructTy(ele0.getContext(), elements, name, isPacked);
   }
 
   template <typename... Args>
   static typename std::enable_if_t<llvm::are_base_of<LLVMType, Args...>::value,
                                    LLVMType>
   createStructTy(StringRef name, LLVMType elt1, Args... elts) {
     SmallVector<LLVMType, 8> fields({elt1, elts...});
     Optional<StringRef> opt_name(name);
-    return createStructTy(&elt1.getDialect(), fields, opt_name);
+    return createStructTy(elt1.getContext(), fields, opt_name);
   }
 
   static LLVMType setStructTyBody(LLVMType structType,
                                   ArrayRef<LLVMType> elements,
                                   bool isPacked = false);
 
   template <typename... Args>
   static typename std::enable_if_t<llvm::are_base_of<LLVMType, Args...>::value,

mlir/lib/Conversion/GPUCommon/ConvertLaunchFuncToRuntimeCalls.cpp

 /// Intermediate data structures are allocated on the stack.
 class GpuLaunchFuncToGpuRuntimeCallsPass
     : public ConvertGpuLaunchFuncToGpuRuntimeCallsBase<
           GpuLaunchFuncToGpuRuntimeCallsPass> {
 private:
   LLVM::LLVMDialect *getLLVMDialect() { return llvmDialect; }
 
   void initializeCachedTypes() {
-    llvmVoidType = LLVM::LLVMType::getVoidTy(llvmDialect);
-    llvmPointerType = LLVM::LLVMType::getInt8PtrTy(llvmDialect);
+    llvmVoidType = LLVM::LLVMType::getVoidTy(&getContext());
+    llvmPointerType = LLVM::LLVMType::getInt8PtrTy(&getContext());
     llvmPointerPointerType = llvmPointerType.getPointerTo();
-    llvmInt8Type = LLVM::LLVMType::getInt8Ty(llvmDialect);
-    llvmInt32Type = LLVM::LLVMType::getInt32Ty(llvmDialect);
-    llvmInt64Type = LLVM::LLVMType::getInt64Ty(llvmDialect);
+    llvmInt8Type = LLVM::LLVMType::getInt8Ty(&getContext());
+    llvmInt32Type = LLVM::LLVMType::getInt32Ty(&getContext());
+    llvmInt64Type = LLVM::LLVMType::getInt64Ty(&getContext());
     llvmIntPtrType = LLVM::LLVMType::getIntNTy(
-        llvmDialect, llvmDialect->getDataLayout().getPointerSizeInBits());
+        &getContext(), llvmDialect->getDataLayout().getPointerSizeInBits());
   }
 
   LLVM::LLVMType getVoidType() { return llvmVoidType; }
 
   LLVM::LLVMType getPointerType() { return llvmPointerType; }
 
   LLVM::LLVMType getPointerPointerType() { return llvmPointerPointerType; }
 
   LLVM::LLVMType getInt8Type() { return llvmInt8Type; }
 
   LLVM::LLVMType getInt32Type() { return llvmInt32Type; }
 
   LLVM::LLVMType getInt64Type() { return llvmInt64Type; }
 
   LLVM::LLVMType getIntPtrType() {
     return LLVM::LLVMType::getIntNTy(
-        getLLVMDialect(),
+        &getContext(),
         getLLVMDialect()->getDataLayout().getPointerSizeInBits());
   }
 
   // Allocate a void pointer on the stack.
   Value allocatePointer(OpBuilder &builder, Location loc) {
     auto one = builder.create<LLVM::ConstantOp>(loc, getInt32Type(),
                                                 builder.getI32IntegerAttr(1));
     return builder.create<LLVM::AllocaOp>(loc, getPointerPointerType(), one,
Value GpuLaunchFuncToGpuRuntimeCallsPass::generateKernelNameConstant(
   // Make sure the trailing zero is included in the constant.
   std::vector<char> kernelName(name.begin(), name.end());
   kernelName.push_back('\0');
 
   std::string globalName =
       std::string(llvm::formatv("{0}_{1}_kernel_name", moduleName, name));
   return LLVM::createGlobalString(
       loc, builder, globalName, StringRef(kernelName.data(), kernelName.size()),
-      LLVM::Linkage::Internal, llvmDialect);
+      LLVM::Linkage::Internal);
 }
 
 // Emits LLVM IR to launch a kernel function. Expects the module that contains
 // the compiled kernel function as a cubin in the 'nvvm.cubin' attribute, or a
 // hsaco in the 'rocdl.hsaco' attribute of the kernel function in the IR.
 //
 // %0 = call %binarygetter
 // %1 = call %moduleLoad(%0)
void GpuLaunchFuncToGpuRuntimeCallsPass::translateGpuLaunchCalls(
   if (!binaryAttr) {
     kernelModule.emitOpError()
         << "missing " << gpuBinaryAnnotation << " attribute";
     return signalPassFailure();
   }
 
   SmallString<128> nameBuffer(kernelModule.getName());
   nameBuffer.append(kGpuBinaryStorageSuffix);
-  Value data = LLVM::createGlobalString(
-      loc, builder, nameBuffer.str(), binaryAttr.getValue(),
-      LLVM::Linkage::Internal, getLLVMDialect());
+  Value data =
+      LLVM::createGlobalString(loc, builder, nameBuffer.str(),
+                               binaryAttr.getValue(), LLVM::Linkage::Internal);
 
   // Emit the load module call to load the module data. Error checking is done
   // in the called helper function.
   auto gpuModuleLoad =
       getOperation().lookupSymbol<LLVM::LLVMFuncOp>(kGpuModuleLoadName);
   auto module = builder.create<LLVM::CallOp>(
       loc, ArrayRef<Type>{getPointerType()},
       builder.getSymbolRefAttr(gpuModuleLoad), ArrayRef<Value>{data});

mlir/lib/Conversion/GPUCommon/GPUOpsLowering.h

auto llvmFuncOp = rewriter.create<LLVM::LLVMFuncOp>(
       // the original function, before the body region is inlined in the new
       // function to maintain the relation between block arguments and the
       // parent operation that assigns their semantics.
       OpBuilder::InsertionGuard guard(rewriter);
 
       // Rewrite workgroup memory attributions to addresses of global buffers.
       rewriter.setInsertionPointToStart(&gpuFuncOp.front());
       unsigned numProperArguments = gpuFuncOp.getNumArguments();
-      auto i32Type = LLVM::LLVMType::getInt32Ty(typeConverter.getDialect());
+      auto i32Type = LLVM::LLVMType::getInt32Ty(rewriter.getContext());
 
       Value zero = nullptr;
       if (!workgroupBuffers.empty())
         zero = rewriter.create<LLVM::ConstantOp>(loc, i32Type,
                                                  rewriter.getI32IntegerAttr(0));
       for (auto en : llvm::enumerate(workgroupBuffers)) {
         LLVM::GlobalOp global = en.value();
         Value address = rewriter.create<LLVM::AddressOfOp>(loc, global);
auto llvmFuncOp = rewriter.create<LLVM::LLVMFuncOp>(
         auto descr = MemRefDescriptor::fromStaticShape(
             rewriter, loc, typeConverter, type, memory);
         signatureConversion.remapInput(numProperArguments + en.index(), descr);
       }
 
       // Rewrite private memory attributions to alloca'ed buffers.
       unsigned numWorkgroupAttributions =
           gpuFuncOp.getNumWorkgroupAttributions();
-      auto int64Ty = LLVM::LLVMType::getInt64Ty(typeConverter.getDialect());
+      auto int64Ty = LLVM::LLVMType::getInt64Ty(rewriter.getContext());
       for (auto en : llvm::enumerate(gpuFuncOp.getPrivateAttributions())) {
         Value attribution = en.value();
         auto type = attribution.getType().cast<MemRefType>();
         assert(type && type.hasStaticShape() &&
                "unexpected type in attribution");
 
         // Explicitly drop memory space when lowering private memory
         // attributions since NVVM models it as `alloca`s in the default

mlir/lib/Conversion/GPUCommon/IndexIntrinsicsOpLowering.h

explicit GPUIndexIntrinsicOpLowering(LLVMTypeConverter &typeConverter)
                              typeConverter),
         indexBitwidth(typeConverter.getIndexTypeBitwidth()) {}
 
   // Convert the kernel arguments to an LLVM type, preserve the rest.
   LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     auto loc = op->getLoc();
-    auto dialect = typeConverter.getDialect();
+    MLIRContext *context = rewriter.getContext();
     Value newOp;
     switch (dimensionToIndex(cast<Op>(op))) {
     case X:
-      newOp = rewriter.create<XOp>(loc, LLVM::LLVMType::getInt32Ty(dialect));
+      newOp = rewriter.create<XOp>(loc, LLVM::LLVMType::getInt32Ty(context));
       break;
     case Y:
-      newOp = rewriter.create<YOp>(loc, LLVM::LLVMType::getInt32Ty(dialect));
+      newOp = rewriter.create<YOp>(loc, LLVM::LLVMType::getInt32Ty(context));
       break;
     case Z:
-      newOp = rewriter.create<ZOp>(loc, LLVM::LLVMType::getInt32Ty(dialect));
+      newOp = rewriter.create<ZOp>(loc, LLVM::LLVMType::getInt32Ty(context));
       break;
     default:
       return failure();
     }
 
     if (indexBitwidth > 32) {
       newOp = rewriter.create<LLVM::SExtOp>(
-          loc, LLVM::LLVMType::getIntNTy(dialect, indexBitwidth), newOp);
+          loc, LLVM::LLVMType::getIntNTy(context, indexBitwidth), newOp);
     } else if (indexBitwidth < 32) {
       newOp = rewriter.create<LLVM::TruncOp>(
-          loc, LLVM::LLVMType::getIntNTy(dialect, indexBitwidth), newOp);
+          loc, LLVM::LLVMType::getIntNTy(context, indexBitwidth), newOp);
     }
 
     rewriter.replaceOp(op, {newOp});
     return success();
   }
 };
 
 } // namespace mlir
 
 #endif // MLIR_CONVERSION_GPUCOMMON_INDEXINTRINSICSOPLOWERING_H_

mlir/lib/Conversion/GPUCommon/OpToFuncCallLowering.h

 
 private:
   Value maybeCast(Value operand, PatternRewriter &rewriter) const {
     LLVM::LLVMType type = operand.getType().cast<LLVM::LLVMType>();
     if (!type.isHalfTy())
       return operand;
 
     return rewriter.create<LLVM::FPExtOp>(
-        operand.getLoc(), LLVM::LLVMType::getFloatTy(&type.getDialect()),
+        operand.getLoc(), LLVM::LLVMType::getFloatTy(rewriter.getContext()),
         operand);
   }
 
   LLVM::LLVMType getFunctionType(LLVM::LLVMType resultType,
                                  ArrayRef<Value> operands) const {
     using LLVM::LLVMType;
     SmallVector<LLVMType, 1> operandTypes;
     for (Value operand : operands) {

mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp

struct GPUShuffleOpLowering : public ConvertToLLVMPattern {
   ///     %shfl_pred = llvm.extractvalue %shfl[1 : index] :
   ///         !llvm<"{ float, i1 }">
   LogicalResult
   matchAndRewrite(Operation *op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     Location loc = op->getLoc();
     gpu::ShuffleOpAdaptor adaptor(operands);
 
-    auto dialect = typeConverter.getDialect();
     auto valueTy = adaptor.value().getType().cast<LLVM::LLVMType>();
-    auto int32Type = LLVM::LLVMType::getInt32Ty(dialect);
-    auto predTy = LLVM::LLVMType::getInt1Ty(dialect);
-    auto resultTy = LLVM::LLVMType::getStructTy(dialect, {valueTy, predTy});
+    auto int32Type = LLVM::LLVMType::getInt32Ty(rewriter.getContext());
+    auto predTy = LLVM::LLVMType::getInt1Ty(rewriter.getContext());
+    auto resultTy =
+        LLVM::LLVMType::getStructTy(rewriter.getContext(), {valueTy, predTy});
 
     Value one = rewriter.create<LLVM::ConstantOp>(
         loc, int32Type, rewriter.getI32IntegerAttr(1));
     // Bit mask of active lanes: `(1 << activeWidth) - 1`.
     Value activeMask = rewriter.create<LLVM::SubOp>(
         loc, int32Type,
         rewriter.create<LLVM::ShlOp>(loc, int32Type, one, adaptor.width()),
         one);

mlir/lib/Conversion/GPUToVulkan/ConvertLaunchFuncToVulkanCalls.cpp

 /// * setNumWorkGroups     -- sets the number of a local workgroups
 /// * runOnVulkan          -- runs vulkan runtime
 /// * deinitVulkan         -- deinitializes vulkan runtime
 ///
 class VulkanLaunchFuncToVulkanCallsPass
     : public ConvertVulkanLaunchFuncToVulkanCallsBase<
           VulkanLaunchFuncToVulkanCallsPass> {
 private:
-  LLVM::LLVMDialect *getLLVMDialect() { return llvmDialect; }
-
   void initializeCachedTypes() {
-    llvmDialect = getContext().getRegisteredDialect<LLVM::LLVMDialect>();
-    llvmFloatType = LLVM::LLVMType::getFloatTy(llvmDialect);
-    llvmVoidType = LLVM::LLVMType::getVoidTy(llvmDialect);
-    llvmPointerType = LLVM::LLVMType::getInt8PtrTy(llvmDialect);
-    llvmInt32Type = LLVM::LLVMType::getInt32Ty(llvmDialect);
-    llvmInt64Type = LLVM::LLVMType::getInt64Ty(llvmDialect);
+    llvmFloatType = LLVM::LLVMType::getFloatTy(&getContext());
+    llvmVoidType = LLVM::LLVMType::getVoidTy(&getContext());
+    llvmPointerType = LLVM::LLVMType::getInt8PtrTy(&getContext());
+    llvmInt32Type = LLVM::LLVMType::getInt32Ty(&getContext());
+    llvmInt64Type = LLVM::LLVMType::getInt64Ty(&getContext());
   }
 
   LLVM::LLVMType getMemRefType(uint32_t rank, LLVM::LLVMType elemenType) {
     // According to the MLIR doc memref argument is converted into a
     // pointer-to-struct argument of type:
     // template <typename Elem, size_t Rank>
     // struct {
     //   Elem *allocated;
     //   Elem *aligned;
     //   int64_t offset;
     //   int64_t sizes[Rank]; // omitted when rank == 0
     //   int64_t strides[Rank]; // omitted when rank == 0
     // };
     auto llvmPtrToElementType = elemenType.getPointerTo();
     auto llvmArrayRankElementSizeType =
         LLVM::LLVMType::getArrayTy(getInt64Type(), rank);
 
     // Create a type
     // `!llvm<"{ `element-type`*, `element-type`*, i64,
     // [`rank` x i64], [`rank` x i64]}">`.
     return LLVM::LLVMType::getStructTy(
-        llvmDialect,
+        &getContext(),
         {llvmPtrToElementType, llvmPtrToElementType, getInt64Type(),
          llvmArrayRankElementSizeType, llvmArrayRankElementSizeType});
   }
 
   LLVM::LLVMType getVoidType() { return llvmVoidType; }
   LLVM::LLVMType getPointerType() { return llvmPointerType; }
   LLVM::LLVMType getInt32Type() { return llvmInt32Type; }
   LLVM::LLVMType getInt64Type() { return llvmInt64Type; }
StringRef stringifyType(LLVM::LLVMType type) {
 
     llvm_unreachable("unsupported type");
   }
 
 public:
   void runOnOperation() override;
 
 private:
-  LLVM::LLVMDialect *llvmDialect;
   LLVM::LLVMType llvmFloatType;
   LLVM::LLVMType llvmVoidType;
   LLVM::LLVMType llvmPointerType;
   LLVM::LLVMType llvmInt32Type;
   LLVM::LLVMType llvmInt64Type;
 
   // TODO: Use an associative array to support multiple vulkan launch calls.
   std::pair<StringAttr, StringAttr> spirvAttributes;
for (auto en :
     }
 
     auto symbolName =
         llvm::formatv("bindMemRef{0}D{1}", rank, stringifyType(type)).str();
     // Special case for fp16 type. Since it is not a supported type in C we use
     // int16_t and bitcast the descriptor.
     if (type.isHalfTy()) {
       auto memRefTy =
-          getMemRefType(rank, LLVM::LLVMType::getInt16Ty(llvmDialect));
+          getMemRefType(rank, LLVM::LLVMType::getInt16Ty(&getContext()));
       ptrToMemRefDescriptor = builder.create<LLVM::BitcastOp>(
           loc, memRefTy.getPointerTo(), ptrToMemRefDescriptor);
     }
     // Create call to `bindMemRef`.
     builder.create<LLVM::CallOp>(
         loc, ArrayRef<Type>{getVoidType()},
         builder.getSymbolRefAttr(
             StringRef(symbolName.data(), symbolName.size())),
void VulkanLaunchFuncToVulkanCallsPass::declareVulkanFunctions(Location loc) {
   if (!module.lookupSymbol(kRunOnVulkan)) {
     builder.create<LLVM::LLVMFuncOp>(
         loc, kRunOnVulkan,
         LLVM::LLVMType::getFunctionTy(getVoidType(), {getPointerType()},
                                       /*isVarArg=*/false));
   }
 
   for (unsigned i = 1; i <= 3; i++) {
-    for (LLVM::LLVMType type : {LLVM::LLVMType::getFloatTy(llvmDialect),
-                                LLVM::LLVMType::getInt32Ty(llvmDialect),
-                                LLVM::LLVMType::getInt16Ty(llvmDialect),
-                                LLVM::LLVMType::getInt8Ty(llvmDialect),
-                                LLVM::LLVMType::getHalfTy(llvmDialect)}) {
+    for (LLVM::LLVMType type : {LLVM::LLVMType::getFloatTy(&getContext()),
+                                LLVM::LLVMType::getInt32Ty(&getContext()),
+                                LLVM::LLVMType::getInt16Ty(&getContext()),
+                                LLVM::LLVMType::getInt8Ty(&getContext()),
+                                LLVM::LLVMType::getHalfTy(&getContext())}) {
       std::string fnName = "bindMemRef" + std::to_string(i) + "D" +
                            std::string(stringifyType(type));
       if (type.isHalfTy())
-        type = getMemRefType(i, LLVM::LLVMType::getInt16Ty(llvmDialect));
+        type = getMemRefType(i, LLVM::LLVMType::getInt16Ty(&getContext()));
       if (!module.lookupSymbol(fnName)) {
         auto fnType = LLVM::LLVMType::getFunctionTy(
             getVoidType(),
             {getPointerType(), getInt32Type(), getInt32Type(),
              getMemRefType(i, type).getPointerTo()},
             /*isVarArg=*/false);
         builder.create<LLVM::LLVMFuncOp>(loc, fnName, fnType);
       }
Value VulkanLaunchFuncToVulkanCallsPass::createEntryPointNameConstant(
     StringRef name, Location loc, OpBuilder &builder) {
   SmallString<16> shaderName(name.begin(), name.end());
   // Append `\0` to follow C style string given that LLVM::createGlobalString()
   // won't handle this directly for us.
   shaderName.push_back('\0');
 
   std::string entryPointGlobalName = (name + "_spv_entry_point_name").str();
   return LLVM::createGlobalString(loc, builder, entryPointGlobalName,
-                                  shaderName, LLVM::Linkage::Internal,
-                                  getLLVMDialect());
+                                  shaderName, LLVM::Linkage::Internal);
 }
 
 void VulkanLaunchFuncToVulkanCallsPass::translateVulkanLaunchCall(
     LLVM::CallOp cInterfaceVulkanLaunchCallOp) {
   OpBuilder builder(cInterfaceVulkanLaunchCallOp);
   Location loc = cInterfaceVulkanLaunchCallOp.getLoc();
   // Create call to `initVulkan`.
   auto initVulkanCall = builder.create<LLVM::CallOp>(
       loc, ArrayRef<Type>{getPointerType()},
       builder.getSymbolRefAttr(kInitVulkan), ArrayRef<Value>{});
   // The result of `initVulkan` function is a pointer to Vulkan runtime, we
   // need to pass that pointer to each Vulkan runtime call.
   auto vulkanRuntime = initVulkanCall.getResult(0);
 
   // Create LLVM global with SPIR-V binary data, so we can pass a pointer with
   // that data to runtime call.
   Value ptrToSPIRVBinary = LLVM::createGlobalString(
       loc, builder, kSPIRVBinary, spirvAttributes.first.getValue(),
-      LLVM::Linkage::Internal, getLLVMDialect());
+      LLVM::Linkage::Internal);
 
   // Create LLVM constant for the size of SPIR-V binary shader.
   Value binarySize = builder.create<LLVM::ConstantOp>(
       loc, getInt32Type(),
       builder.getI32IntegerAttr(spirvAttributes.first.getValue().size()));
 
   // Create call to `bindMemRef` for each memref operand.
   createBindMemRefCalls(cInterfaceVulkanLaunchCallOp, vulkanRuntime);

mlir/lib/Conversion/SPIRVToLLVM/ConvertSPIRVToLLVM.cpp

 
 /// Converts SPIR-V struct with no offset to packed LLVM struct.
 static Type convertStructTypePacked(spirv::StructType type,
                                     LLVMTypeConverter &converter) {
   auto elementsVector = llvm::to_vector<8>(
       llvm::map_range(type.getElementTypes(), [&](Type elementType) {
         return converter.convertType(elementType).cast<LLVM::LLVMType>();
       }));
-  return LLVM::LLVMType::getStructTy(converter.getDialect(), elementsVector,
+  return LLVM::LLVMType::getStructTy(type.getContext(), elementsVector,
                                      /*isPacked=*/true);
 }
 
 /// Creates LLVM dialect constant with the given value.
 static Value createI32ConstantOf(Location loc, PatternRewriter &rewriter,
-                                 LLVMTypeConverter &converter, unsigned value) {
+                                 unsigned value) {
   return rewriter.create<LLVM::ConstantOp>(
-      loc, LLVM::LLVMType::getInt32Ty(converter.getDialect()),
+      loc, LLVM::LLVMType::getInt32Ty(rewriter.getContext()),
       rewriter.getIntegerAttr(rewriter.getI32Type(), value));
 }
 
 /// Utility for `spv.Load` and `spv.Store` conversion.
 static LogicalResult replaceWithLoadOrStore(Operation *op,
                                             ConversionPatternRewriter &rewriter,
                                             LLVMTypeConverter &typeConverter,
                                             unsigned alignment, bool isVolatile,
matchAndRewrite(spirv::VariableOp varOp, ArrayRef<Value> operands,
     if (init && !pointerTo.isIntOrFloat() && !pointerTo.isa<VectorType>())
       return failure();
 
     auto dstType = typeConverter.convertType(srcType);
     if (!dstType)
       return failure();
 
     Location loc = varOp.getLoc();
-    Value size = createI32ConstantOf(loc, rewriter, typeConverter, 1);
+    Value size = createI32ConstantOf(loc, rewriter, 1);
     if (!init) {
       rewriter.replaceOpWithNewOp<LLVM::AllocaOp>(varOp, dstType, size);
       return success();
     }
     Value allocated = rewriter.create<LLVM::AllocaOp>(loc, dstType, size);
     rewriter.create<LLVM::StoreOp>(loc, init, allocated);
     rewriter.replaceOp(varOp, allocated);
     return success();

mlir/lib/Conversion/StandardToLLVM/StandardToLLVM.cpp

 }
 
 /// Get the LLVM context.
 llvm::LLVMContext &LLVMTypeConverter::getLLVMContext() {
   return llvmDialect->getLLVMContext();
 }
 
 LLVM::LLVMType LLVMTypeConverter::getIndexType() {
-  return LLVM::LLVMType::getIntNTy(llvmDialect, getIndexTypeBitwidth());
+  return LLVM::LLVMType::getIntNTy(&getContext(), getIndexTypeBitwidth());
 }
 
 unsigned LLVMTypeConverter::getPointerBitwidth(unsigned addressSpace) {
   return llvmDialect->getDataLayout().getPointerSizeInBits(addressSpace);
 }
 
 Type LLVMTypeConverter::convertIndexType(IndexType type) {
   return getIndexType();
 }
 
 Type LLVMTypeConverter::convertIntegerType(IntegerType type) {
-  return LLVM::LLVMType::getIntNTy(llvmDialect, type.getWidth());
+  return LLVM::LLVMType::getIntNTy(&getContext(), type.getWidth());
 }
 
 Type LLVMTypeConverter::convertFloatType(FloatType type) {
   switch (type.getKind()) {
   case mlir::StandardTypes::F32:
-    return LLVM::LLVMType::getFloatTy(llvmDialect);
+    return LLVM::LLVMType::getFloatTy(&getContext());
   case mlir::StandardTypes::F64:
-    return LLVM::LLVMType::getDoubleTy(llvmDialect);
+    return LLVM::LLVMType::getDoubleTy(&getContext());
   case mlir::StandardTypes::F16:
-    return LLVM::LLVMType::getHalfTy(llvmDialect);
+    return LLVM::LLVMType::getHalfTy(&getContext());
   case mlir::StandardTypes::BF16: {
-    return LLVM::LLVMType::getBFloatTy(llvmDialect);
+    return LLVM::LLVMType::getBFloatTy(&getContext());
   }
   default:
     llvm_unreachable("non-float type in convertFloatType");
   }
 }
 
 // Convert a `ComplexType` to an LLVM type. The result is a complex number
 // struct with entries for the
 //   1. real part and for the
 //   2. imaginary part.
 static constexpr unsigned kRealPosInComplexNumberStruct = 0;
 static constexpr unsigned kImaginaryPosInComplexNumberStruct = 1;
 Type LLVMTypeConverter::convertComplexType(ComplexType type) {
   auto elementType = convertType(type.getElementType()).cast<LLVM::LLVMType>();
-  return LLVM::LLVMType::getStructTy(llvmDialect, {elementType, elementType});
+  return LLVM::LLVMType::getStructTy(&getContext(), {elementType, elementType});
 }
 
 // Except for signatures, MLIR function types are converted into LLVM
 // pointer-to-function types.
 Type LLVMTypeConverter::convertFunctionType(FunctionType type) {
   SignatureConversion conversion(type.getNumInputs());
   LLVM::LLVMType converted =
       convertFunctionSignature(type, /*isVariadic=*/false, conversion);
LLVMTypeConverter::convertMemRefSignature(MemRefType type) {
   auto rank = type.getRank();
   results.insert(results.end(), 2 * rank, indexTy);
   return results;
 }
 
 /// In signatures, unranked MemRef descriptors are expanded into a pair "rank,
 /// pointer to descriptor".
 SmallVector<Type, 2> LLVMTypeConverter::convertUnrankedMemRefSignature() {
-  return {getIndexType(), LLVM::LLVMType::getInt8PtrTy(llvmDialect)};
+  return {getIndexType(), LLVM::LLVMType::getInt8PtrTy(&getContext())};
 }
 
 // Function types are converted to LLVM Function types by recursively converting
 // argument and result types.  If MLIR Function has zero results, the LLVM
 // Function has one VoidType result.  If MLIR Function has more than one result,
 // they are into an LLVM StructType in their order of appearance.
 LLVM::LLVMType LLVMTypeConverter::convertFunctionSignature(
     FunctionType type, bool isVariadic,
LLVM::LLVMType LLVMTypeConverter::convertFunctionSignature(
   for (Type type : result.getConvertedTypes())
     argTypes.push_back(unwrap(type));
 
   // If function does not return anything, create the void result type,
   // if it returns on element, convert it, otherwise pack the result types into
   // a struct.
   LLVM::LLVMType resultType =
       type.getNumResults() == 0
-          ? LLVM::LLVMType::getVoidTy(llvmDialect)
+          ? LLVM::LLVMType::getVoidTy(&getContext())
           : unwrap(packFunctionResults(type.getResults()));
   if (!resultType)
     return {};
   return LLVM::LLVMType::getFunctionTy(resultType, argTypes, isVariadic);
 }
 
 /// Converts the function type to a C-compatible format, in particular using
 /// pointers to memref descriptors for arguments.
 LLVM::LLVMType
 LLVMTypeConverter::convertFunctionTypeCWrapper(FunctionType type) {
   SmallVector<LLVM::LLVMType, 4> inputs;
 
   for (Type t : type.getInputs()) {
     auto converted = convertType(t).dyn_cast_or_null<LLVM::LLVMType>();
     if (!converted)
       return {};
     if (t.isa<MemRefType, UnrankedMemRefType>())
       converted = converted.getPointerTo();
     inputs.push_back(converted);
   }
 
   LLVM::LLVMType resultType =
       type.getNumResults() == 0
-          ? LLVM::LLVMType::getVoidTy(llvmDialect)
+          ? LLVM::LLVMType::getVoidTy(&getContext())
           : unwrap(packFunctionResults(type.getResults()));
   if (!resultType)
     return {};
 
   return LLVM::LLVMType::getFunctionTy(resultType, inputs, false);
 }
 
 // Convert a MemRef to an LLVM type. The result is a MemRef descriptor which

 //    stack allocated (alloca) copy of a MemRef descriptor that got casted to
 //    be unranked.
 
 static constexpr unsigned kRankInUnrankedMemRefDescriptor = 0;
 static constexpr unsigned kPtrInUnrankedMemRefDescriptor = 1;
 
 Type LLVMTypeConverter::convertUnrankedMemRefType(UnrankedMemRefType type) {
   auto rankTy = getIndexType();
-  auto ptrTy = LLVM::LLVMType::getInt8PtrTy(llvmDialect);
+  auto ptrTy = LLVM::LLVMType::getInt8PtrTy(&getContext());
   return LLVM::LLVMType::getStructTy(rankTy, ptrTy);
 }
 
 // Convert an n-D vector type to an LLVM vector type via (n-1)-D array type when
 // n > 1.
 // For example, `vector<4 x f32>` converts to `!llvm.type<"<4 x float>">` and
 // `vector<4 x 8 x 16 f32>` converts to `!llvm<"[4 x [8 x <16 x float>]]">`.
 Type LLVMTypeConverter::convertVectorType(VectorType type) {
llvm::LLVMContext &ConvertToLLVMPattern::getContext() const {
   return typeConverter.getLLVMContext();
 }
 
 LLVM::LLVMType ConvertToLLVMPattern::getIndexType() const {
   return typeConverter.getIndexType();
 }
 
 LLVM::LLVMType ConvertToLLVMPattern::getVoidType() const {
-  return LLVM::LLVMType::getVoidTy(&getDialect());
+  return LLVM::LLVMType::getVoidTy(&typeConverter.getContext());
 }
 
 LLVM::LLVMType ConvertToLLVMPattern::getVoidPtrType() const {
-  return LLVM::LLVMType::getInt8PtrTy(&getDialect());
+  return LLVM::LLVMType::getInt8PtrTy(&typeConverter.getContext());
 }
 
 Value ConvertToLLVMPattern::createIndexConstant(
     ConversionPatternRewriter &builder, Location loc, uint64_t value) const {
   return createIndexAttrConstant(builder, loc, getIndexType(), value);
 }
 
 Value ConvertToLLVMPattern::linearizeSubscripts(
if (unrankedMemrefs.empty())
     return success();
 
   // Compute allocation sizes.
   SmallVector<Value, 4> sizes;
   UnrankedMemRefDescriptor::computeSizes(builder, loc, typeConverter,
                                          unrankedMemrefs, sizes);
 
   // Get frequently used types.
-  auto voidType = LLVM::LLVMType::getVoidTy(typeConverter.getDialect());
-  auto voidPtrType = LLVM::LLVMType::getInt8PtrTy(typeConverter.getDialect());
-  auto i1Type = LLVM::LLVMType::getInt1Ty(typeConverter.getDialect());
+  MLIRContext *context = builder.getContext();
+  auto voidType = LLVM::LLVMType::getVoidTy(context);
+  auto voidPtrType = LLVM::LLVMType::getInt8PtrTy(context);
+  auto i1Type = LLVM::LLVMType::getInt1Ty(context);
   LLVM::LLVMType indexType = typeConverter.getIndexType();
 
   // Find the malloc and free, or declare them if necessary.
   auto module = builder.getInsertionPoint()->getParentOfType<ModuleOp>();
   auto mallocFunc = module.lookupSymbol<LLVM::LLVMFuncOp>("malloc");
   if (!mallocFunc && toDynamic) {
     OpBuilder::InsertionGuard guard(builder);
     builder.setInsertionPointToStart(module.getBody());
for (auto &nestedOp : entryBlock.without_terminator()) {
       mapping.map(nestedOp.getResults(), clone->getResults());
     }
     Value result = mapping.lookup(entryBlock.getTerminator()->getOperand(0));
 
     // Prepare the epilog of the loop block.
     // Append the cmpxchg op to the end of the loop block.
     auto successOrdering = LLVM::AtomicOrdering::acq_rel;
     auto failureOrdering = LLVM::AtomicOrdering::monotonic;
-    auto boolType = LLVM::LLVMType::getInt1Ty(&getDialect());
+    auto boolType = LLVM::LLVMType::getInt1Ty(rewriter.getContext());
     auto pairType = LLVM::LLVMType::getStructTy(valueType, boolType);
     auto cmpxchg = rewriter.create<LLVM::AtomicCmpXchgOp>(
         loc, pairType, dataPtr, loopArgument, result, successOrdering,
         failureOrdering);
     // Extract the %new_loaded and %ok values from the pair.
     Value newLoaded = rewriter.create<LLVM::ExtractValueOp>(
         loc, valueType, cmpxchg, rewriter.getI64ArrayAttr({0}));
     Value ok = rewriter.create<LLVM::ExtractValueOp>(
Type LLVMTypeConverter::packFunctionResults(ArrayRef<Type> types) {
   resultTypes.reserve(types.size());
   for (auto t : types) {
     auto converted = convertType(t).dyn_cast<LLVM::LLVMType>();
     if (!converted)
       return {};
     resultTypes.push_back(converted);
   }
 
-  return LLVM::LLVMType::getStructTy(llvmDialect, resultTypes);
+  return LLVM::LLVMType::getStructTy(&getContext(), resultTypes);
 }
 
 Value LLVMTypeConverter::promoteOneMemRefDescriptor(Location loc, Value operand,
                                                     OpBuilder &builder) {
   auto *context = builder.getContext();
-  auto int64Ty = LLVM::LLVMType::getInt64Ty(getDialect());
+  auto int64Ty = LLVM::LLVMType::getInt64Ty(builder.getContext());
   auto indexType = IndexType::get(context);
   // Alloca with proper alignment. We do not expect optimizations of this
   // alloca op and so we omit allocating at the entry block.
   auto ptrType = operand.getType().cast<LLVM::LLVMType>().getPointerTo();
   Value one = builder.create<LLVM::ConstantOp>(loc, int64Ty,
                                                IntegerAttr::get(indexType, 1));
   Value allocated =
       builder.create<LLVM::AllocaOp>(loc, ptrType, one, /*alignment=*/0);

mlir/lib/Conversion/VectorToLLVM/ConvertVectorToLLVM.cpp

if (positionAttrs.size() > 1) {
           ArrayAttr::get(positionAttrs.drop_back(), context);
       extracted = rewriter.create<LLVM::ExtractValueOp>(
           loc, typeConverter.convertType(oneDVectorType), extracted,
           nMinusOnePositionAttrs);
     }
 
     // Remaining extraction of element from 1-D LLVM vector
     auto position = positionAttrs.back().cast<IntegerAttr>();
-    auto i64Type = LLVM::LLVMType::getInt64Ty(typeConverter.getDialect());
+    auto i64Type = LLVM::LLVMType::getInt64Ty(rewriter.getContext());
     auto constant = rewriter.create<LLVM::ConstantOp>(loc, i64Type, position);
     extracted =
         rewriter.create<LLVM::ExtractElementOp>(loc, extracted, constant);
     rewriter.replaceOp(op, extracted);
 
     return success();
   }
 };
if (positionAttrs.size() > 1) {
       auto nMinusOnePositionAttrs =
           ArrayAttr::get(positionAttrs.drop_back(), context);
       extracted = rewriter.create<LLVM::ExtractValueOp>(
           loc, typeConverter.convertType(oneDVectorType), extracted,
           nMinusOnePositionAttrs);
     }
 
     // Insertion of an element into a 1-D LLVM vector.
-    auto i64Type = LLVM::LLVMType::getInt64Ty(typeConverter.getDialect());
+    auto i64Type = LLVM::LLVMType::getInt64Ty(rewriter.getContext());
     auto constant = rewriter.create<LLVM::ConstantOp>(loc, i64Type, position);
     Value inserted = rewriter.create<LLVM::InsertElementOp>(
         loc, typeConverter.convertType(oneDVectorType), extracted,
         adaptor.source(), constant);
 
     // Potential insertion of resulting 1-D vector into array.
     if (positionAttrs.size() > 1) {
       auto nMinusOnePositionAttrs =
if (isContiguous) {
           break;
         }
       }
     }
     // Only contiguous source tensors supported atm.
     if (failed(successStrides) || !isContiguous)
       return failure();
 
-    auto int64Ty = LLVM::LLVMType::getInt64Ty(typeConverter.getDialect());
+    auto int64Ty = LLVM::LLVMType::getInt64Ty(rewriter.getContext());
 
     // Create descriptor.
     auto desc = MemRefDescriptor::undef(rewriter, loc, llvmTargetDescriptorTy);
     Type llvmTargetElementTy = desc.getElementType();
     // Set allocated ptr.
     Value allocated = sourceMemRef.allocatedPtr(rewriter, loc);
     allocated =
         rewriter.create<LLVM::BitcastOp>(loc, llvmTargetElementTy, allocated);
public:
   }
 
 private:
   void emitRanks(ConversionPatternRewriter &rewriter, Operation *op,
                  Value value, VectorType vectorType, Operation *printer,
                  int64_t rank) const {
     Location loc = op->getLoc();
     if (rank == 0) {
-      if (value.getType() ==
-          LLVM::LLVMType::getInt1Ty(typeConverter.getDialect())) {
+      if (value.getType() == LLVM::LLVMType::getInt1Ty(rewriter.getContext())) {
         // Convert i1 (bool) to i32 so we can use the print_i32 method.
         // This avoids the need for a print_i1 method with an unclear ABI.
-        auto i32Type = LLVM::LLVMType::getInt32Ty(typeConverter.getDialect());
+        auto i32Type = LLVM::LLVMType::getInt32Ty(rewriter.getContext());
         auto trueVal = rewriter.create<ConstantOp>(
             loc, i32Type, rewriter.getI32IntegerAttr(1));
         auto falseVal = rewriter.create<ConstantOp>(
             loc, i32Type, rewriter.getI32IntegerAttr(0));
         value = rewriter.create<SelectOp>(loc, value, trueVal, falseVal);
       }
       emitCall(rewriter, loc, printer, value);
       return;
private:
   // Helper to emit a call.
   static void emitCall(ConversionPatternRewriter &rewriter, Location loc,
                        Operation *ref, ValueRange params = ValueRange()) {
     rewriter.create<LLVM::CallOp>(loc, ArrayRef<Type>{},
                                   rewriter.getSymbolRefAttr(ref), params);
   }
 
   // Helper for printer method declaration (first hit) and lookup.
-  static Operation *getPrint(Operation *op, LLVM::LLVMDialect *dialect,
-                             StringRef name, ArrayRef<LLVM::LLVMType> params) {
+  static Operation *getPrint(Operation *op, StringRef name,
+                             ArrayRef<LLVM::LLVMType> params) {
     auto module = op->getParentOfType<ModuleOp>();
     auto func = module.lookupSymbol<LLVM::LLVMFuncOp>(name);
     if (func)
       return func;
     OpBuilder moduleBuilder(module.getBodyRegion());
     return moduleBuilder.create<LLVM::LLVMFuncOp>(
         op->getLoc(), name,
-        LLVM::LLVMType::getFunctionTy(LLVM::LLVMType::getVoidTy(dialect),
-                                      params, /*isVarArg=*/false));
+        LLVM::LLVMType::getFunctionTy(
+            LLVM::LLVMType::getVoidTy(op->getContext()), params,
+            /*isVarArg=*/false));
   }
 
   // Helpers for method names.
   Operation *getPrintI32(Operation *op) const {
-    LLVM::LLVMDialect *dialect = typeConverter.getDialect();
-    return getPrint(op, dialect, "print_i32",
-                    LLVM::LLVMType::getInt32Ty(dialect));
+    return getPrint(op, "print_i32",
+                    LLVM::LLVMType::getInt32Ty(op->getContext()));
   }
   Operation *getPrintI64(Operation *op) const {
-    LLVM::LLVMDialect *dialect = typeConverter.getDialect();
-    return getPrint(op, dialect, "print_i64",
-                    LLVM::LLVMType::getInt64Ty(dialect));
+    return getPrint(op, "print_i64",
+                    LLVM::LLVMType::getInt64Ty(op->getContext()));
   }
   Operation *getPrintFloat(Operation *op) const {
-    LLVM::LLVMDialect *dialect = typeConverter.getDialect();
-    return getPrint(op, dialect, "print_f32",
-                    LLVM::LLVMType::getFloatTy(dialect));
+    return getPrint(op, "print_f32",
+                    LLVM::LLVMType::getFloatTy(op->getContext()));
   }
   Operation *getPrintDouble(Operation *op) const {
-    LLVM::LLVMDialect *dialect = typeConverter.getDialect();
-    return getPrint(op, dialect, "print_f64",
-                    LLVM::LLVMType::getDoubleTy(dialect));
+    return getPrint(op, "print_f64",
+                    LLVM::LLVMType::getDoubleTy(op->getContext()));
   }
   Operation *getPrintOpen(Operation *op) const {
-    return getPrint(op, typeConverter.getDialect(), "print_open", {});
+    return getPrint(op, "print_open", {});
   }
   Operation *getPrintClose(Operation *op) const {
-    return getPrint(op, typeConverter.getDialect(), "print_close", {});
+    return getPrint(op, "print_close", {});
   }
   Operation *getPrintComma(Operation *op) const {
-    return getPrint(op, typeConverter.getDialect(), "print_comma", {});
+    return getPrint(op, "print_comma", {});
   }
   Operation *getPrintNewline(Operation *op) const {
-    return getPrint(op, typeConverter.getDialect(), "print_newline", {});
+    return getPrint(op, "print_newline", {});
   }
 };
 
 /// Progressive lowering of ExtractStridedSliceOp to either:
 ///   1. extractelement + insertelement for the 1-D case
 ///   2. extract + optional strided_slice + insert for the n-D case.
 class VectorStridedSliceOpConversion
     : public OpRewritePattern<ExtractStridedSliceOp> {

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

if (std::is_same<CmpPredicateType, ICmpPredicate>()) {
     predicateValue = static_cast<int64_t>(predicate.getValue());
   }
 
   result.attributes.set("predicate",
                         parser.getBuilder().getI64IntegerAttr(predicateValue));
 
   // The result type is either i1 or a vector type <? x i1> if the inputs are
   // vectors.
-  auto *dialect = builder.getContext()->getRegisteredDialect<LLVMDialect>();
-  auto resultType = LLVMType::getInt1Ty(dialect);
+  auto resultType = LLVMType::getInt1Ty(builder.getContext());
   auto argType = type.dyn_cast<LLVM::LLVMType>();
   if (!argType)
     return parser.emitError(trailingTypeLoc, "expected LLVM IR dialect type");
   if (argType.isVectorTy())
     resultType =
         LLVMType::getVectorTy(resultType, argType.getVectorNumElements());
 
   result.addTypes({resultType});
if (isDirect) {
     result.addTypes(funcType.getResults());
   } else {
     // Construct the LLVM IR Dialect function type that the first operand
     // should match.
     if (funcType.getNumResults() > 1)
       return parser.emitError(trailingTypeLoc,
                               "expected function with 0 or 1 result");
 
-    auto *llvmDialect =
-        builder.getContext()->getRegisteredDialect<LLVM::LLVMDialect>();
     LLVM::LLVMType llvmResultType;
     if (funcType.getNumResults() == 0) {
-      llvmResultType = LLVM::LLVMType::getVoidTy(llvmDialect);
+      llvmResultType = LLVM::LLVMType::getVoidTy(builder.getContext());
     } else {
       llvmResultType = funcType.getResult(0).dyn_cast<LLVM::LLVMType>();
       if (!llvmResultType)
         return parser.emitError(trailingTypeLoc,
                                 "expected result to have LLVM type");
     }
 
     SmallVector<LLVM::LLVMType, 8> argTypes;
static ParseResult parseCallOp(OpAsmParser &parser, OperationState &result) {
   } else {
     // Construct the LLVM IR Dialect function type that the first operand
     // should match.
     if (funcType.getNumResults() > 1)
       return parser.emitError(trailingTypeLoc,
                               "expected function with 0 or 1 result");
 
     Builder &builder = parser.getBuilder();
-    auto *llvmDialect =
-        builder.getContext()->getRegisteredDialect<LLVM::LLVMDialect>();
     LLVM::LLVMType llvmResultType;
     if (funcType.getNumResults() == 0) {
-      llvmResultType = LLVM::LLVMType::getVoidTy(llvmDialect);
+      llvmResultType = LLVM::LLVMType::getVoidTy(builder.getContext());
     } else {
       llvmResultType = funcType.getResult(0).dyn_cast<LLVM::LLVMType>();
       if (!llvmResultType)
         return parser.emitError(trailingTypeLoc,
                                 "expected result to have LLVM type");
     }
 
     SmallVector<LLVM::LLVMType, 8> argTypes;
static ParseResult parseGlobalOp(OpAsmParser &parser, OperationState &result) {
 
   if (types.size() > 1)
     return parser.emitError(parser.getNameLoc(), "expected zero or one type");
 
   Region &initRegion = *result.addRegion();
   if (types.empty()) {
     if (auto strAttr = value.dyn_cast_or_null<StringAttr>()) {
       MLIRContext *context = parser.getBuilder().getContext();
-      auto *dialect = context->getRegisteredDialect<LLVMDialect>();
       auto arrayType = LLVM::LLVMType::getArrayTy(
-          LLVM::LLVMType::getInt8Ty(dialect), strAttr.getValue().size());
+          LLVM::LLVMType::getInt8Ty(context), strAttr.getValue().size());
       types.push_back(arrayType);
     } else {
       return parser.emitError(parser.getNameLoc(),
                               "type can only be omitted for string globals");
     }
   } else if (parser.parseOptionalRegion(initRegion, /*arguments=*/{},
                                         /*argTypes=*/{})) {
     return failure();
for (auto t : inputs) {
     if (!llvmTy) {
       parser.emitError(loc, "failed to construct function type: expected LLVM "
                             "type for function arguments");
       return {};
     }
     llvmInputs.push_back(llvmTy);
   }
 
-  // Get the dialect from the input type, if any exist.  Look it up in the
-  // context otherwise.
-  LLVMDialect *dialect =
-      llvmInputs.empty() ? b.getContext()->getRegisteredDialect<LLVMDialect>()
-                         : &llvmInputs.front().getDialect();
-
   // No output is denoted as "void" in LLVM type system.
-  LLVMType llvmOutput = outputs.empty() ? LLVMType::getVoidTy(dialect)
+  LLVMType llvmOutput = outputs.empty() ? LLVMType::getVoidTy(b.getContext())
                                         : outputs.front().dyn_cast<LLVMType>();
   if (!llvmOutput) {
     parser.emitError(loc, "failed to construct function type: expected LLVM "
                           "type for function results");
     return {};
   }
   return LLVMType::getFunctionTy(llvmOutput, llvmInputs,
                                  variadicFlag.isVariadic());
if (parser.parseOperand(ptr) || parser.parseComma() ||
       parseAtomicOrdering(parser, result, "failure_ordering") ||
       parser.parseOptionalAttrDict(result.attributes) ||
       parser.parseColonType(type) ||
       parser.resolveOperand(ptr, type.getPointerTo(), result.operands) ||
       parser.resolveOperand(cmp, type, result.operands) ||
       parser.resolveOperand(val, type, result.operands))
     return failure();
 
-  auto *dialect = builder.getContext()->getRegisteredDialect<LLVMDialect>();
-  auto boolType = LLVMType::getInt1Ty(dialect);
+  auto boolType = LLVMType::getInt1Ty(builder.getContext());
   auto resultType = LLVMType::getStructTy(type, boolType);
   result.addTypes(resultType);
 
   return success();
 }
 
 static LogicalResult verify(AtomicCmpXchgOp op) {
   auto ptrType = op.ptr().getType().cast<LLVM::LLVMType>();

 }
 
 //===----------------------------------------------------------------------===//
 // Utility functions.
 //===----------------------------------------------------------------------===//
 
 Value mlir::LLVM::createGlobalString(Location loc, OpBuilder &builder,
                                      StringRef name, StringRef value,
-                                     LLVM::Linkage linkage,
-                                     LLVM::LLVMDialect *llvmDialect) {
+                                     LLVM::Linkage linkage) {
   assert(builder.getInsertionBlock() &&
          builder.getInsertionBlock()->getParentOp() &&
          "expected builder to point to a block constrained in an op");
   auto module =
       builder.getInsertionBlock()->getParentOp()->getParentOfType<ModuleOp>();
   assert(module && "builder points to an op outside of a module");
 
   // Create the global at the entry of the module.
   OpBuilder moduleBuilder(module.getBodyRegion());
-  auto type = LLVM::LLVMType::getArrayTy(LLVM::LLVMType::getInt8Ty(llvmDialect),
-                                         value.size());
+  MLIRContext *ctx = builder.getContext();
+  auto type =
+      LLVM::LLVMType::getArrayTy(LLVM::LLVMType::getInt8Ty(ctx), value.size());
   auto global = moduleBuilder.create<LLVM::GlobalOp>(
       loc, type, /*isConstant=*/true, linkage, name,
       builder.getStringAttr(value));
 
   // Get the pointer to the first character in the global string.
   Value globalPtr = builder.create<LLVM::AddressOfOp>(loc, global);
   Value cst0 = builder.create<LLVM::ConstantOp>(
-      loc, LLVM::LLVMType::getInt64Ty(llvmDialect),
+      loc, LLVM::LLVMType::getInt64Ty(ctx),
       builder.getIntegerAttr(builder.getIndexType(), 0));
-  return builder.create<LLVM::GEPOp>(loc,
-                                     LLVM::LLVMType::getInt8PtrTy(llvmDialect),
+  return builder.create<LLVM::GEPOp>(loc, LLVM::LLVMType::getInt8PtrTy(ctx),
                                      globalPtr, ArrayRef<Value>({cst0, cst0}));
 }
 
 bool mlir::LLVM::satisfiesLLVMModule(Operation *op) {
   return op->hasTrait<OpTrait::SymbolTable>() &&
          op->hasTrait<OpTrait::IsIsolatedFromAbove>();
 }
 

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

unsigned LLVMType::getStructNumElements() {
   return cast<LLVMStructType>().getBody().size();
 }
 
 bool LLVMType::isStructTy() { return isa<LLVMStructType>(); }
 
 //----------------------------------------------------------------------------//
 // Utilities used to generate floating point types.
 
-LLVMType LLVMType::getDoubleTy(LLVMDialect *dialect) {
-  return LLVMDoubleType::get(dialect->getContext());
+LLVMType LLVMType::getDoubleTy(MLIRContext *context) {
+  return LLVMDoubleType::get(context);
 }
 
-LLVMType LLVMType::getFloatTy(LLVMDialect *dialect) {
-  return LLVMFloatType::get(dialect->getContext());
+LLVMType LLVMType::getFloatTy(MLIRContext *context) {
+  return LLVMFloatType::get(context);
 }
 
-LLVMType LLVMType::getBFloatTy(LLVMDialect *dialect) {
-  return LLVMBFloatType::get(dialect->getContext());
+LLVMType LLVMType::getBFloatTy(MLIRContext *context) {
+  return LLVMBFloatType::get(context);
 }
 
-LLVMType LLVMType::getHalfTy(LLVMDialect *dialect) {
-  return LLVMHalfType::get(dialect->getContext());
+LLVMType LLVMType::getHalfTy(MLIRContext *context) {
+  return LLVMHalfType::get(context);
 }
 
-LLVMType LLVMType::getFP128Ty(LLVMDialect *dialect) {
-  return LLVMFP128Type::get(dialect->getContext());
+LLVMType LLVMType::getFP128Ty(MLIRContext *context) {
+  return LLVMFP128Type::get(context);
 }
 
-LLVMType LLVMType::getX86_FP80Ty(LLVMDialect *dialect) {
-  return LLVMX86FP80Type::get(dialect->getContext());
+LLVMType LLVMType::getX86_FP80Ty(MLIRContext *context) {
+  return LLVMX86FP80Type::get(context);
 }
 
 //----------------------------------------------------------------------------//
 // Utilities used to generate integer types.
 
-LLVMType LLVMType::getIntNTy(LLVMDialect *dialect, unsigned numBits) {
-  return LLVMIntegerType::get(dialect->getContext(), numBits);
+LLVMType LLVMType::getIntNTy(MLIRContext *context, unsigned numBits) {
+  return LLVMIntegerType::get(context, numBits);
 }
 
 //----------------------------------------------------------------------------//
 // Utilities used to generate other miscellaneous types.
 
 LLVMType LLVMType::getArrayTy(LLVMType elementType, uint64_t numElements) {
   return LLVMArrayType::get(elementType, numElements);
 }
 
 LLVMType LLVMType::getFunctionTy(LLVMType result, ArrayRef<LLVMType> params,
                                  bool isVarArg) {
   return LLVMFunctionType::get(result, params, isVarArg);
 }
 
-LLVMType LLVMType::getStructTy(LLVMDialect *dialect,
+LLVMType LLVMType::getStructTy(MLIRContext *context,
                                ArrayRef<LLVMType> elements, bool isPacked) {
-  return LLVMStructType::getLiteral(dialect->getContext(), elements, isPacked);
+  return LLVMStructType::getLiteral(context, elements, isPacked);
 }
 
 LLVMType LLVMType::getVectorTy(LLVMType elementType, unsigned numElements) {
   return LLVMFixedVectorType::get(elementType, numElements);
 }
 
 //----------------------------------------------------------------------------//
 // Void type utilities.
 
-LLVMType LLVMType::getVoidTy(LLVMDialect *dialect) {
-  return LLVMVoidType::get(dialect->getContext());
+LLVMType LLVMType::getVoidTy(MLIRContext *context) {
+  return LLVMVoidType::get(context);
 }
 
 bool LLVMType::isVoidTy() { return isa<LLVMVoidType>(); }
 
 //----------------------------------------------------------------------------//
 // Creation and setting of LLVM's identified struct types
 
-LLVMType LLVMType::createStructTy(LLVMDialect *dialect,
+LLVMType LLVMType::createStructTy(MLIRContext *context,
                                   ArrayRef<LLVMType> elements,
                                   Optional<StringRef> name, bool isPacked) {
   assert(name.hasValue() &&
          "identified structs with no identifier not supported");
   StringRef stringNameBase = name.getValueOr("");
   std::string stringName = stringNameBase.str();
   unsigned counter = 0;
   do {
-    auto type =
-        LLVMStructType::getIdentified(dialect->getContext(), stringName);
+    auto type = LLVMStructType::getIdentified(context, stringName);
     if (type.isInitialized() || failed(type.setBody(elements, isPacked))) {
       counter += 1;
       stringName =
           (Twine(stringNameBase) + "." + std::to_string(counter)).str();
       continue;
     }
     return type;
   } while (true);

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

 //===----------------------------------------------------------------------===//
 
 static void printNVVMIntrinsicOp(OpAsmPrinter &p, Operation *op) {
   p << op->getName() << " " << op->getOperands();
   if (op->getNumResults() > 0)
     p << " : " << op->getResultTypes();
 }
 
-static LLVM::LLVMDialect *getLlvmDialect(OpAsmParser &parser) {
-  return parser.getBuilder()
-      .getContext()
-      ->getRegisteredDialect<LLVM::LLVMDialect>();
-}
-
 // <operation> ::=
 //     `llvm.nvvm.shfl.sync.bfly %dst, %val, %offset, %clamp_and_mask`
 //      ({return_value_and_is_valid})? : result_type
 static ParseResult parseNVVMShflSyncBflyOp(OpAsmParser &parser,
                                            OperationState &result) {
   SmallVector<OpAsmParser::OperandType, 8> ops;
   Type resultType;
   if (parser.parseOperandList(ops) ||
       parser.parseOptionalAttrDict(result.attributes) ||
       parser.parseColonType(resultType) ||
       parser.addTypeToList(resultType, result.types))
     return failure();
 
   auto type = resultType.cast<LLVM::LLVMType>();
   for (auto &attr : result.attributes) {
     if (attr.first != "return_value_and_is_valid")
       continue;
     if (type.isStructTy() && type.getStructNumElements() > 0)
       type = type.getStructElementType(0);
     break;
   }
 
-  auto int32Ty = LLVM::LLVMType::getInt32Ty(getLlvmDialect(parser));
+  auto int32Ty = LLVM::LLVMType::getInt32Ty(parser.getBuilder().getContext());
   return parser.resolveOperands(ops, {int32Ty, type, int32Ty, int32Ty},
                                 parser.getNameLoc(), result.operands);
 }
 
 // <operation> ::= `llvm.nvvm.vote.ballot.sync %mask, %pred` : result_type
 static ParseResult parseNVVMVoteBallotOp(OpAsmParser &parser,
                                          OperationState &result) {
-  auto llvmDialect = getLlvmDialect(parser);
-  auto int32Ty = LLVM::LLVMType::getInt32Ty(llvmDialect);
-  auto int1Ty = LLVM::LLVMType::getInt1Ty(llvmDialect);
+  MLIRContext *context = parser.getBuilder().getContext();
+  auto int32Ty = LLVM::LLVMType::getInt32Ty(context);
+  auto int1Ty = LLVM::LLVMType::getInt1Ty(context);
 
   SmallVector<OpAsmParser::OperandType, 8> ops;
   Type type;
   return failure(parser.parseOperandList(ops) ||
                  parser.parseOptionalAttrDict(result.attributes) ||
                  parser.parseColonType(type) ||
                  parser.addTypeToList(type, result.types) ||
                  parser.resolveOperands(ops, {int32Ty, int1Ty},
                                         parser.getNameLoc(), result.operands));
 }
 
 static LogicalResult verify(MmaOp op) {
-  auto dialect = op.getContext()->getRegisteredDialect<LLVM::LLVMDialect>();
-  auto f16Ty = LLVM::LLVMType::getHalfTy(dialect);
+  MLIRContext *context = op.getContext();
+  auto f16Ty = LLVM::LLVMType::getHalfTy(context);
   auto f16x2Ty = LLVM::LLVMType::getVectorTy(f16Ty, 2);
-  auto f32Ty = LLVM::LLVMType::getFloatTy(dialect);
+  auto f32Ty = LLVM::LLVMType::getFloatTy(context);
   auto f16x2x4StructTy = LLVM::LLVMType::getStructTy(
-      dialect, {f16x2Ty, f16x2Ty, f16x2Ty, f16x2Ty});
+      context, {f16x2Ty, f16x2Ty, f16x2Ty, f16x2Ty});
   auto f32x8StructTy = LLVM::LLVMType::getStructTy(
-      dialect, {f32Ty, f32Ty, f32Ty, f32Ty, f32Ty, f32Ty, f32Ty, f32Ty});
+      context, {f32Ty, f32Ty, f32Ty, f32Ty, f32Ty, f32Ty, f32Ty, f32Ty});
 
   SmallVector<Type, 12> operand_types(op.getOperandTypes().begin(),
                                       op.getOperandTypes().end());
   if (operand_types != SmallVector<Type, 8>(8, f16x2Ty) &&
       operand_types != SmallVector<Type, 12>{f16x2Ty, f16x2Ty, f16x2Ty, f16x2Ty,
                                              f32Ty, f32Ty, f32Ty, f32Ty, f32Ty,
                                              f32Ty, f32Ty, f32Ty}) {
     return op.emitOpError(

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

 
 using namespace mlir;
 using namespace ROCDL;
 
 //===----------------------------------------------------------------------===//
 // Parsing for ROCDL ops
 //===----------------------------------------------------------------------===//
 
-static LLVM::LLVMDialect *getLlvmDialect(OpAsmParser &parser) {
-  return parser.getBuilder()
-      .getContext()
-      ->getRegisteredDialect<LLVM::LLVMDialect>();
-}
-
 // <operation> ::=
 //     `llvm.amdgcn.buffer.load.* %rsrc, %vindex, %offset, %glc, %slc :
 //     result_type`
 static ParseResult parseROCDLMubufLoadOp(OpAsmParser &parser,
                                          OperationState &result) {
   SmallVector<OpAsmParser::OperandType, 8> ops;
   Type type;
   if (parser.parseOperandList(ops, 5) || parser.parseColonType(type) ||
       parser.addTypeToList(type, result.types))
     return failure();
 
-  auto int32Ty = LLVM::LLVMType::getInt32Ty(getLlvmDialect(parser));
-  auto int1Ty = LLVM::LLVMType::getInt1Ty(getLlvmDialect(parser));
+  MLIRContext *context = parser.getBuilder().getContext();
+  auto int32Ty = LLVM::LLVMType::getInt32Ty(context);
+  auto int1Ty = LLVM::LLVMType::getInt1Ty(context);
   auto i32x4Ty = LLVM::LLVMType::getVectorTy(int32Ty, 4);
   return parser.resolveOperands(ops,
                                 {i32x4Ty, int32Ty, int32Ty, int1Ty, int1Ty},
                                 parser.getNameLoc(), result.operands);
 }
 
 // <operation> ::=
 //     `llvm.amdgcn.buffer.store.* %vdata, %rsrc, %vindex, %offset, %glc, %slc :
 //     result_type`
 static ParseResult parseROCDLMubufStoreOp(OpAsmParser &parser,
                                           OperationState &result) {
   SmallVector<OpAsmParser::OperandType, 8> ops;
   Type type;
   if (parser.parseOperandList(ops, 6) || parser.parseColonType(type))
     return failure();
 
-  auto int32Ty = LLVM::LLVMType::getInt32Ty(getLlvmDialect(parser));
-  auto int1Ty = LLVM::LLVMType::getInt1Ty(getLlvmDialect(parser));
+  MLIRContext *context = parser.getBuilder().getContext();
+  auto int32Ty = LLVM::LLVMType::getInt32Ty(context);
+  auto int1Ty = LLVM::LLVMType::getInt1Ty(context);
   auto i32x4Ty = LLVM::LLVMType::getVectorTy(int32Ty, 4);
 
   if (parser.resolveOperands(ops,
                              {type, i32x4Ty, int32Ty, int32Ty, int1Ty, int1Ty},
                              parser.getNameLoc(), result.operands))
     return failure();
   return success();
 }