[mlir] Enable opaque pointers in LLVM conversion passes by default

[zero9178]

Summary

Part of https://discourse.llvm.org/t/rfc-switching-the-llvm-dialect-and-dialect-lowerings-to-opaque-pointers/68179

When this patch lands any downstream users with custom LLVM conversion passes not yet using opaque pointers will start either experiencing assertions being triggered, null pointer dereferences or at the very least verifier errors. These can be either fixed by switching to opaque pointers or simply disabling opaque pointers in both pass options of any upstream conversion passes and any uses of LLVMTypeConverter via the LowerToLLVMOptions.

Users using just MLIRs conversion passes to the LLVM Dialect should not experience any change in functionality except when inspecting the output from the passes.

---

As was outlined in https://discourse.llvm.org/t/psa-in-tree-conversion-passes-can-now-be-used-with-llvm-opaque-pointers-please-switch-your-downstream-projects/68738, I'd land this patch on Friday with an accompanying PSA as that is exactly two weeks after the initial anncouncement.

[gysit]

LGTM!

flang/lib/Optimizer/CodeGen/TypeConverter.h

 
 namespace fir {
 
 /// FIR type converter
 /// This converts FIR types to LLVM types (for now)
 class LLVMTypeConverter : public mlir::LLVMTypeConverter {
 public:
   LLVMTypeConverter(mlir::ModuleOp module, bool applyTBAA)
-      : mlir::LLVMTypeConverter(module.getContext()),
+      : mlir::LLVMTypeConverter(module.getContext(),
+                                [&] {
+                                  mlir::LowerToLLVMOptions options(
+                                      module.getContext());
+                                  options.useOpaquePointers = false;
+                                  return options;
+                                }()),
         kindMapping(getKindMapping(module)),
         specifics(CodeGenSpecifics::get(module.getContext(),
                                         getTargetTriple(module),
                                         getKindMapping(module))),
         tbaaBuilder(module, applyTBAA) {
     LLVM_DEBUG(llvm::dbgs() << "FIR type converter\n");
 
     // Each conversion should return a value of type mlir::Type.

mlir/include/mlir/Conversion/LLVMCommon/LoweringOptions.h

 /// Options to control the LLVM lowering. The struct is used to share lowering
 /// options between passes, patterns, and type converter.
 class LowerToLLVMOptions {
 public:
   explicit LowerToLLVMOptions(MLIRContext *ctx);
   LowerToLLVMOptions(MLIRContext *ctx, const DataLayout &dl);
 
   bool useBarePtrCallConv = false;
-  bool useOpaquePointers = false;
+  bool useOpaquePointers = true;
 
   enum class AllocLowering {
     /// Use malloc for for heap allocations.
     Malloc,
 
     /// Use aligned_alloc for heap allocations.
     AlignedAlloc,
 

mlir/include/mlir/Conversion/Passes.td

def ConvertAsyncToLLVMPass : Pass<"convert-async-to-llvm", "ModuleOp"> {
   let dependentDialects = [
     "arith::ArithDialect",
     "async::AsyncDialect",
     "LLVM::LLVMDialect",
     "func::FuncDialect",
   ];
   let options = [
     Option<"useOpaquePointers", "use-opaque-pointers", "bool",
-           /*default=*/"false", "Generate LLVM IR using opaque pointers "
+           /*default=*/"true", "Generate LLVM IR using opaque pointers "
            "instead of typed pointers">,
   ];
 }
 
 //===----------------------------------------------------------------------===//
 // BufferizationToMemRef
 //===----------------------------------------------------------------------===//
 
let description = [{
     already present in the IR will be kept as is.
   }];
   let dependentDialects = ["LLVM::LLVMDialect"];
   let options = [
     Option<"indexBitwidth", "index-bitwidth", "unsigned",
            /*default=kDeriveIndexBitwidthFromDataLayout*/"0",
            "Bitwidth of the index type, 0 to use size of machine word">,
     Option<"useOpaquePointers", "use-opaque-pointers", "bool",
-                   /*default=*/"false", "Generate LLVM IR using opaque pointers "
+                   /*default=*/"true", "Generate LLVM IR using opaque pointers "
                    "instead of typed pointers">,
   ];
 }
 
 //===----------------------------------------------------------------------===//
 // ControlFlowToSPIRV
 //===----------------------------------------------------------------------===//
 
let options = [
     Option<"indexBitwidth", "index-bitwidth", "unsigned",
            /*default=kDeriveIndexBitwidthFromDataLayout*/"0",
            "Bitwidth of the index type, 0 to use size of machine word">,
     Option<"dataLayout", "data-layout", "std::string",
            /*default=*/"\"\"",
            "String description (LLVM format) of the data layout that is "
            "expected on the produced module">,
     Option<"useOpaquePointers", "use-opaque-pointers", "bool",
-                       /*default=*/"false", "Generate LLVM IR using opaque pointers "
+                       /*default=*/"true", "Generate LLVM IR using opaque pointers "
                        "instead of typed pointers">,
   ];
 }
 
 //===----------------------------------------------------------------------===//
 // FuncToSPIRV
 //===----------------------------------------------------------------------===//
 
Option<"kernelBarePtrCallConv", "use-bare-pointers-for-kernels", "bool",
              "Use bare pointers to pass memref arguments to kernels. "
              "The kernel must use the same setting for this option."
            >,
     Option<"gpuBinaryAnnotation", "gpu-binary-annotation", "std::string",
                /*default=*/"gpu::getDefaultGpuBinaryAnnotation()",
                "Annotation attribute string for GPU binary"
                >,
     Option<"useOpaquePointers", "use-opaque-pointers", "bool",
-               /*default=*/"false", "Generate LLVM IR using opaque pointers "
+               /*default=*/"true", "Generate LLVM IR using opaque pointers "
                "instead of typed pointers">,
   ];
 
   let dependentDialects = [
     "LLVM::LLVMDialect",
     "memref::MemRefDialect",
   ];
 }
let description = [{
     linked to the variables in the kernel module, and a series of copies to/from
     them to emulate the memory transfer from the host or to the device sides. It
     also converts the remaining Arithmetic, Func, and MemRef dialects into LLVM
     dialect, emitting C wrappers.
   }];
 
   let options = [
     Option<"useOpaquePointers", "use-opaque-pointers", "bool",
-                 /*default=*/"false", "Generate LLVM IR using opaque pointers "
+                 /*default=*/"true", "Generate LLVM IR using opaque pointers "
                  "instead of typed pointers">
   ];
 
   let dependentDialects = ["LLVM::LLVMDialect"];
 }
 
 //===----------------------------------------------------------------------===//
 // GPUToNVVM
def ConvertGpuOpsToNVVMOps : Pass<"convert-gpu-to-nvvm", "gpu::GPUModuleOp"> {
   ];
   let options = [
     Option<"indexBitwidth", "index-bitwidth", "unsigned",
            /*default=kDeriveIndexBitwidthFromDataLayout*/"0",
            "Bitwidth of the index type, 0 to use size of machine word">,
     Option<"hasRedux", "has-redux", "bool", /*default=*/"false",
            "Target gpu supports redux">,
     Option<"useOpaquePointers", "use-opaque-pointers", "bool",
-                   /*default=*/"false", "Generate LLVM IR using opaque pointers "
+                   /*default=*/"true", "Generate LLVM IR using opaque pointers "
                    "instead of typed pointers">,
   ];
 }
 
 //===----------------------------------------------------------------------===//
 // GPUToROCDL
 //===----------------------------------------------------------------------===//
 
Option<"runtime", "runtime", "::mlir::gpu::amd::Runtime",
           "::mlir::gpu::amd::Runtime::Unknown",
           "Runtime code will be run on (default is Unknown, can also use HIP or OpenCl)",
           [{::llvm::cl::values(
             clEnumValN(::mlir::gpu::amd::Runtime::Unknown, "unknown", "Unknown (default)"),
             clEnumValN(::mlir::gpu::amd::Runtime::HIP, "HIP", "HIP"),
             clEnumValN(::mlir::gpu::amd::Runtime::OpenCL, "OpenCL", "OpenCL")
           )}]>,
     Option<"useOpaquePointers", "use-opaque-pointers", "bool",
-               /*default=*/"false", "Generate LLVM IR using opaque pointers "
+               /*default=*/"true", "Generate LLVM IR using opaque pointers "
                "instead of typed pointers">,
   ];
 }
 
 //===----------------------------------------------------------------------===//
 // GPUToSPIRV
 //===----------------------------------------------------------------------===//
 
def ConvertVulkanLaunchFuncToVulkanCallsPass
   let summary = "Convert vulkanLaunch external call to Vulkan runtime external "
                 "calls";
   let description = [{
     This pass is only intended for the mlir-vulkan-runner.
   }];
 
   let options = [
      Option<"useOpaquePointers", "use-opaque-pointers", "bool",
-            /*default=*/"false", "Generate LLVM IR using opaque pointers "
+            /*default=*/"true", "Generate LLVM IR using opaque pointers "
             "instead of typed pointers">
   ];
 
   let dependentDialects = ["LLVM::LLVMDialect"];
 }
 
 //===----------------------------------------------------------------------===//
 // ConvertIndexToLLVMPass

 //===----------------------------------------------------------------------===//
 // LinalgToLLVM
 //===----------------------------------------------------------------------===//
 
 def ConvertLinalgToLLVMPass : Pass<"convert-linalg-to-llvm", "ModuleOp"> {
   let summary = "Convert the operations from the linalg dialect into the LLVM "
                 "dialect";
   let dependentDialects = ["scf::SCFDialect", "LLVM::LLVMDialect"];
-
   let options = [
      Option<"useOpaquePointers", "use-opaque-pointers", "bool",
-            /*default=*/"false", "Generate LLVM IR using opaque pointers "
+                /*default=*/"true", "Generate LLVM IR using opaque pointers "
                 "instead of typed pointers">
   ];
 }
 
 //===----------------------------------------------------------------------===//
 // LinalgToStandard
 //===----------------------------------------------------------------------===//
 
 def ConvertLinalgToStandard : Pass<"convert-linalg-to-std", "ModuleOp"> {
Option<"indexBitwidth", "index-bitwidth", "unsigned",
            /*default=kDeriveIndexBitwidthFromDataLayout*/"0",
            "Bitwidth of the index type, 0 to use size of machine word">,
     Option<"useGenericFunctions", "use-generic-functions",
            "bool",
            /*default=*/"false",
            "Use generic allocation and deallocation functions instead of the "
            "classic 'malloc', 'aligned_alloc' and 'free' functions">,
     Option<"useOpaquePointers", "use-opaque-pointers", "bool",
-               /*default=*/"false", "Generate LLVM IR using opaque pointers "
+               /*default=*/"true", "Generate LLVM IR using opaque pointers "
                "instead of typed pointers">
   ];
 }
 
 //===----------------------------------------------------------------------===//
 // MemRefToSPIRV
 //===----------------------------------------------------------------------===//
 
let description = [{
     This pass converts supported NVGPU ops to NVVM dialect intrinsics.
   }];
 
   let dependentDialects = [
     "NVVM::NVVMDialect",
   ];
   let options = [
     Option<"useOpaquePointers", "use-opaque-pointers", "bool",
-              /*default=*/"false", "Generate LLVM IR using opaque pointers "
+              /*default=*/"true", "Generate LLVM IR using opaque pointers "
               "instead of typed pointers">
   ];
 }
 
 
 //===----------------------------------------------------------------------===//
 // OpenACCToSCF
 //===----------------------------------------------------------------------===//
 
 def ConvertOpenACCToSCF : Pass<"convert-openacc-to-scf", "ModuleOp"> {
   let summary = "Convert the OpenACC ops to OpenACC with SCF dialect";
   let constructor = "mlir::createConvertOpenACCToSCFPass()";
   let dependentDialects = ["scf::SCFDialect", "acc::OpenACCDialect"];
 }
 
 //===----------------------------------------------------------------------===//
 // OpenACCToLLVM
 //===----------------------------------------------------------------------===//
 
 def ConvertOpenACCToLLVMPass : Pass<"convert-openacc-to-llvm", "ModuleOp"> {
   let summary = "Convert the OpenACC ops to LLVM dialect";
   let dependentDialects = ["LLVM::LLVMDialect"];
   let options = [
     Option<"useOpaquePointers", "use-opaque-pointers", "bool",
-           /*default=*/"false", "Generate LLVM IR using opaque pointers "
+           /*default=*/"true", "Generate LLVM IR using opaque pointers "
            "instead of typed pointers">,
   ];
 }
 
 //===----------------------------------------------------------------------===//
 // OpenMPToLLVM
 //===----------------------------------------------------------------------===//
 

 //===----------------------------------------------------------------------===//
 
 def ConvertSCFToOpenMPPass : Pass<"convert-scf-to-openmp", "ModuleOp"> {
   let summary = "Convert SCF parallel loop to OpenMP parallel + workshare "
                 "constructs.";
 
   let options = [
     Option<"useOpaquePointers", "use-opaque-pointers", "bool",
-                 /*default=*/"false", "Generate LLVM IR using opaque pointers "
+                 /*default=*/"true", "Generate LLVM IR using opaque pointers "
                  "instead of typed pointers">
   ];
 
   let dependentDialects = ["omp::OpenMPDialect", "LLVM::LLVMDialect",
                            "memref::MemRefDialect"];
 }
 
 //===----------------------------------------------------------------------===//
def ConvertSPIRVToLLVMPass : Pass<"convert-spirv-to-llvm", "ModuleOp"> {
   let description = [{
     See https://mlir.llvm.org/docs/SPIRVToLLVMDialectConversion/
     for more details.
   }];
   let dependentDialects = ["LLVM::LLVMDialect"];
 
   let options = [
     Option<"useOpaquePointers", "use-opaque-pointers", "bool",
-                 /*default=*/"false", "Generate LLVM IR using opaque pointers "
+                 /*default=*/"true", "Generate LLVM IR using opaque pointers "
                  "instead of typed pointers">
   ];
 }
 
 //===----------------------------------------------------------------------===//
 // TensorToLinalg
 //===----------------------------------------------------------------------===//
 
Option<"armSVE", "enable-arm-sve",
            "bool", /*default=*/"false",
            "Enables the use of ArmSVE dialect while lowering the vector "
        "dialect.">,
     Option<"x86Vector", "enable-x86vector",
            "bool", /*default=*/"false",
            "Enables the use of X86Vector dialect while lowering the vector "
      "dialect.">,
     Option<"useOpaquePointers", "use-opaque-pointers", "bool",
-               /*default=*/"false", "Generate LLVM IR using opaque pointers "
+               /*default=*/"true", "Generate LLVM IR using opaque pointers "
                "instead of typed pointers">
   ];
 }
 
 //===----------------------------------------------------------------------===//
 // VectorToSPIRV
 //===----------------------------------------------------------------------===//
 
 def ConvertVectorToSPIRV : Pass<"convert-vector-to-spirv"> {
   let summary = "Convert Vector dialect to SPIR-V dialect";
   let constructor = "mlir::createConvertVectorToSPIRVPass()";
   let dependentDialects = ["spirv::SPIRVDialect"];
 }
 
 #endif // MLIR_CONVERSION_PASSES

mlir/test/mlir-cpu-runner/bare-ptr-call-conv.mlir

scf.for %arg2 = %c0 to %c2 step %c1 {
     %3 = arith.addf %1, %cst_0 : f32
     memref.store %3, %arg1[%arg2] : memref<2xf32>
     // CHECK-NEXT: 4, 4
   }
   return
 }
 
 // External declarations.
-llvm.func @malloc(i64) -> !llvm.ptr<i8>
-llvm.func @free(!llvm.ptr<i8>)
 func.func private @printF32(%arg0: f32)
 func.func private @printComma()
 func.func private @printNewline()
 
 func.func @main()
 {
   %c2 = arith.constant 2 : index
   %c0 = arith.constant 0 : index