[mlir][SPIRVToLLVM] Add pass option to emit opaque-pointers

[zero9178]

Summary

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

This patch adds the pass option and required changes to the patterns to support the emission of LLVM IR opaque pointers. Given how close SPIRV semantics are to LLVM IR semantics this boils down to just a few changes:

• Making sure that GEP and alloca are built with the explicit base pointer type
• creating opaque pointers instead of typed pointers if requested
• omitting pointer to pointer bitcasts

[antiagainst]

Nice, thanks!

mlir/include/mlir/Conversion/Passes.td

let description = [{
 
     This transformation creates a sequence of global variables that are later
     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 "
+                 "instead of typed pointers">
+  ];
+
   let dependentDialects = ["LLVM::LLVMDialect"];
 }
 
 //===----------------------------------------------------------------------===//
 // GPUToNVVM
 //===----------------------------------------------------------------------===//
 
 def ConvertGpuOpsToNVVMOps : Pass<"convert-gpu-to-nvvm", "gpu::GPUModuleOp"> {

 
 def ConvertSPIRVToLLVMPass : Pass<"convert-spirv-to-llvm", "ModuleOp"> {
   let summary = "Convert SPIR-V dialect to LLVM dialect";
   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 "
+                 "instead of typed pointers">
+  ];
 }
 
 //===----------------------------------------------------------------------===//
 // TensorToLinalg
 //===----------------------------------------------------------------------===//
 
 def ConvertTensorToLinalg : Pass<"convert-tensor-to-linalg", "ModuleOp"> {
   let summary = "Convert some Tensor dialect ops to Linalg dialect";

mlir/lib/Conversion/SPIRVToLLVM/ConvertLaunchFuncToLLVMCalls.cpp

for (const auto &operand : llvm::enumerate(kernelOperands)) {
             /*isConstant=*/false, LLVM::Linkage::Linkonce, name, Attribute(),
             /*alignment=*/0);
         rewriter.setInsertionPoint(launchOp);
       }
 
       // Copy the data from src operand pointer to dst global variable. Save
       // src, dst and size so that we can copy data back after emulating the
       // kernel call.
-      Value dst = rewriter.create<LLVM::AddressOfOp>(loc, dstGlobal);
+      Value dst = rewriter.create<LLVM::AddressOfOp>(
+          loc, typeConverter->convertType(spirvGlobal.getType()),
+          dstGlobal.getSymName());
       copy(loc, dst, src, sizeBytes, rewriter);
 
       CopyInfo info;
       info.dst = dst;
       info.src = src;
       info.size = sizeBytes;
       copyInfo.push_back(info);
     }
void runOnOperation() override {
     // Request C wrapper emission.
     for (auto func : module.getOps<func::FuncOp>()) {
       func->setAttr(LLVM::LLVMDialect::getEmitCWrapperAttrName(),
                     UnitAttr::get(&getContext()));
     }
 
     // Specify options to lower to LLVM and pull in the conversion patterns.
     LowerToLLVMOptions options(module.getContext());
+    options.useOpaquePointers = useOpaquePointers;
+
     auto *context = module.getContext();
     RewritePatternSet patterns(context);
     LLVMTypeConverter typeConverter(context, options);
     mlir::arith::populateArithToLLVMConversionPatterns(typeConverter, patterns);
     populateFinalizeMemRefToLLVMConversionPatterns(typeConverter, patterns);
     populateFuncToLLVMConversionPatterns(typeConverter, patterns);
     patterns.add<GPULaunchLowering>(typeConverter);
 

mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVM.cpp

static std::optional<Type> convertArrayType(spirv::ArrayType type,
   auto llvmElementType = converter.convertType(elementType);
   unsigned numElements = type.getNumElements();
   return LLVM::LLVMArrayType::get(llvmElementType, numElements);
 }
 
 /// Converts SPIR-V pointer type to LLVM pointer. Pointer's storage class is not
 /// modelled at the moment.
 static Type convertPointerType(spirv::PointerType type,
-                               TypeConverter &converter) {
+                               LLVMTypeConverter &converter) {
   auto pointeeType = converter.convertType(type.getPointeeType());
-  return LLVM::LLVMPointerType::get(pointeeType);
+  return converter.getPointerType(pointeeType);
 }
 
 /// Converts SPIR-V runtime array to LLVM array. Since LLVM allows indexing over
 /// the bounds, the runtime array is converted to a 0-sized LLVM array. There is
 /// no modelling of array stride at the moment.
 static std::optional<Type> convertRuntimeArrayType(spirv::RuntimeArrayType type,
                                                    TypeConverter &converter) {
   if (type.getArrayStride() != 0)
matchAndRewrite(spirv::AccessChainOp op, OpAdaptor adaptor,
     auto indices = llvm::to_vector<4>(adaptor.getIndices());
     Type indexType = op.getIndices().front().getType();
     auto llvmIndexType = typeConverter.convertType(indexType);
     if (!llvmIndexType)
       return failure();
     Value zero = rewriter.create<LLVM::ConstantOp>(
         op.getLoc(), llvmIndexType, rewriter.getIntegerAttr(indexType, 0));
     indices.insert(indices.begin(), zero);
-    rewriter.replaceOpWithNewOp<LLVM::GEPOp>(op, dstType, adaptor.getBasePtr(),
-                                             indices);
+    rewriter.replaceOpWithNewOp<LLVM::GEPOp>(
+        op, dstType,
+        typeConverter.convertType(op.getBasePtr()
+                                      .getType()
+                                      .cast<spirv::PointerType>()
+                                      .getPointeeType()),
+        adaptor.getBasePtr(), indices);
     return success();
   }
 };
 
 class AddressOfPattern : public SPIRVToLLVMConversion<spirv::AddressOfOp> {
 public:
   using SPIRVToLLVMConversion<spirv::AddressOfOp>::SPIRVToLLVMConversion;
 
matchAndRewrite(spirv::VariableOp varOp, OpAdaptor adaptor,
 
     auto dstType = typeConverter.convertType(srcType);
     if (!dstType)
       return failure();
 
     Location loc = varOp.getLoc();
     Value size = createI32ConstantOf(loc, rewriter, 1);
     if (!init) {
-      rewriter.replaceOpWithNewOp<LLVM::AllocaOp>(varOp, dstType, size);
+      rewriter.replaceOpWithNewOp<LLVM::AllocaOp>(
+          varOp, dstType, typeConverter.convertType(pointerTo), size);
       return success();
     }
-    Value allocated = rewriter.create<LLVM::AllocaOp>(loc, dstType, size);
+    Value allocated = rewriter.create<LLVM::AllocaOp>(
+        loc, dstType, typeConverter.convertType(pointerTo), size);
     rewriter.create<LLVM::StoreOp>(loc, adaptor.getInitializer(), allocated);
     rewriter.replaceOp(varOp, allocated);
     return success();
   }
 };
 
 //===----------------------------------------------------------------------===//
+// BitcastOp conversion
+//===----------------------------------------------------------------------===//
+
+class BitcastConversionPattern
+    : public SPIRVToLLVMConversion<spirv::BitcastOp> {
+public:
+  using SPIRVToLLVMConversion<spirv::BitcastOp>::SPIRVToLLVMConversion;
+
+  LogicalResult
+  matchAndRewrite(spirv::BitcastOp bitcastOp, OpAdaptor adaptor,
+                  ConversionPatternRewriter &rewriter) const override {
+    auto dstType = typeConverter.convertType(bitcastOp.getType());
+    if (!dstType)
+      return failure();
+
+    if (typeConverter.useOpaquePointers() &&
+        dstType.isa<LLVM::LLVMPointerType>()) {
+      rewriter.replaceOp(bitcastOp, adaptor.getOperand());
+      return success();
+    }
+
+    rewriter.replaceOpWithNewOp<LLVM::BitcastOp>(
+        bitcastOp, dstType, adaptor.getOperands(), bitcastOp->getAttrs());
+    return success();
+  }
+};
+
+//===----------------------------------------------------------------------===//
 // FuncOp conversion
 //===----------------------------------------------------------------------===//
 
 class FuncConversionPattern : public SPIRVToLLVMConversion<spirv::FuncOp> {
 public:
   using SPIRVToLLVMConversion<spirv::FuncOp>::SPIRVToLLVMConversion;
 
   LogicalResult
patterns.add<
       DirectConversionPattern<spirv::BitCountOp, LLVM::CtPopOp>,
       DirectConversionPattern<spirv::BitReverseOp, LLVM::BitReverseOp>,
       DirectConversionPattern<spirv::BitwiseAndOp, LLVM::AndOp>,
       DirectConversionPattern<spirv::BitwiseOrOp, LLVM::OrOp>,
       DirectConversionPattern<spirv::BitwiseXorOp, LLVM::XOrOp>,
       NotPattern<spirv::NotOp>,
 
       // Cast ops
-      DirectConversionPattern<spirv::BitcastOp, LLVM::BitcastOp>,
+      BitcastConversionPattern,
       DirectConversionPattern<spirv::ConvertFToSOp, LLVM::FPToSIOp>,
       DirectConversionPattern<spirv::ConvertFToUOp, LLVM::FPToUIOp>,
       DirectConversionPattern<spirv::ConvertSToFOp, LLVM::SIToFPOp>,
       DirectConversionPattern<spirv::ConvertUToFOp, LLVM::UIToFPOp>,
       IndirectCastPattern<spirv::FConvertOp, LLVM::FPExtOp, LLVM::FPTruncOp>,
       IndirectCastPattern<spirv::SConvertOp, LLVM::SExtOp, LLVM::TruncOp>,
       IndirectCastPattern<spirv::UConvertOp, LLVM::ZExtOp, LLVM::TruncOp>,
 

mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.cpp

 public:
   using Base::Base;
 };
 } // namespace
 
 void ConvertSPIRVToLLVMPass::runOnOperation() {
   MLIRContext *context = &getContext();
   ModuleOp module = getOperation();
-  LLVMTypeConverter converter(&getContext());
+
+  LowerToLLVMOptions options(&getContext());
+  options.useOpaquePointers = useOpaquePointers;
+
+  LLVMTypeConverter converter(&getContext(), options);
 
   // Encode global variable's descriptor set and binding if they exist.
   encodeBindAttribute(module);
 
   RewritePatternSet patterns(context);
 
   populateSPIRVToLLVMTypeConversion(converter);
 

mlir/test/Conversion/SPIRVToLLVM/arithmetic-ops-to-llvm.mlir

-// RUN: mlir-opt -convert-spirv-to-llvm %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // spirv.IAdd
 //===----------------------------------------------------------------------===//
 
 // CHECK-LABEL: @iadd_scalar
 spirv.func @iadd_scalar(%arg0: i32, %arg1: i32) "None" {
   // CHECK: llvm.add %{{.*}}, %{{.*}} : i32

mlir/test/Conversion/SPIRVToLLVM/bitwise-ops-to-llvm.mlir

-// RUN: mlir-opt -convert-spirv-to-llvm %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // spirv.BitCount
 //===----------------------------------------------------------------------===//
 
 // CHECK-LABEL: @bitcount_scalar
 spirv.func @bitcount_scalar(%arg0: i16) "None" {
   // CHECK: llvm.intr.ctpop(%{{.*}}) : (i16) -> i16

mlir/test/Conversion/SPIRVToLLVM/cast-ops-to-llvm.mlir

-// RUN: mlir-opt -convert-spirv-to-llvm %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // spirv.Bitcast
 //===----------------------------------------------------------------------===//
 
 // CHECK-LABEL: @bitcast_float_to_integer_scalar
 spirv.func @bitcast_float_to_integer_scalar(%arg0 : f32) "None" {
   // CHECK: llvm.bitcast {{.*}} : f32 to i32

 spirv.func @bitcast_vector_to_vector(%arg0 : vector<4xf32>) "None" {
   // CHECK: {{.*}} = llvm.bitcast {{.*}} : vector<4xf32> to vector<2xi64>
   %0 = spirv.Bitcast %arg0: vector<4xf32> to vector<2xi64>
   spirv.Return
 }
 
 // CHECK-LABEL: @bitcast_pointer
 spirv.func @bitcast_pointer(%arg0: !spirv.ptr<f32, Function>) "None" {
-  // CHECK: llvm.bitcast %{{.*}} : !llvm.ptr<f32> to !llvm.ptr<i32>
+  // CHECK-NOT: llvm.bitcast
   %0 = spirv.Bitcast %arg0 : !spirv.ptr<f32, Function> to !spirv.ptr<i32, Function>
   spirv.Return
 }
 
 //===----------------------------------------------------------------------===//
 // spirv.ConvertFToS
 //===----------------------------------------------------------------------===//
 

mlir/test/Conversion/SPIRVToLLVM/comparison-ops-to-llvm.mlir

-// RUN: mlir-opt -convert-spirv-to-llvm %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // spirv.IEqual
 //===----------------------------------------------------------------------===//
 
 // CHECK-LABEL: @i_equal_scalar
 spirv.func @i_equal_scalar(%arg0: i32, %arg1: i32) "None" {
   // CHECK: llvm.icmp "eq" %{{.*}}, %{{.*}} : i32

mlir/test/Conversion/SPIRVToLLVM/constant-op-to-llvm.mlir

-// RUN: mlir-opt -convert-spirv-to-llvm %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // spirv.Constant
 //===----------------------------------------------------------------------===//
 
 // CHECK-LABEL: @bool_constant_scalar
 spirv.func @bool_constant_scalar() "None" {
   // CHECK: llvm.mlir.constant(true) : i1

mlir/test/Conversion/SPIRVToLLVM/control-flow-ops-to-llvm.mlir

-// RUN: mlir-opt -convert-spirv-to-llvm -split-input-file -verify-diagnostics %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' -split-input-file -verify-diagnostics %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // spirv.Branch
 //===----------------------------------------------------------------------===//
 
 spirv.module Logical GLSL450 {
   spirv.func @branch_without_arguments() -> () "None" {
     // CHECK: llvm.br ^bb1

mlir/test/Conversion/SPIRVToLLVM/func-ops-to-llvm.mlir

-// RUN: mlir-opt -convert-spirv-to-llvm %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // spirv.Return
 //===----------------------------------------------------------------------===//
 
 // CHECK-LABEL: @return
 spirv.func @return() "None" {
   // CHECK: llvm.return

mlir/test/Conversion/SPIRVToLLVM/gl-ops-to-llvm.mlir

-// RUN: mlir-opt -convert-spirv-to-llvm %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // spirv.GL.Ceil
 //===----------------------------------------------------------------------===//
 
 // CHECK-LABEL: @ceil
 spirv.func @ceil(%arg0: f32, %arg1: vector<3xf16>) "None" {
   // CHECK: llvm.intr.ceil(%{{.*}}) : (f32) -> f32

mlir/test/Conversion/SPIRVToLLVM/logical-ops-to-llvm.mlir

-// RUN: mlir-opt -convert-spirv-to-llvm %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // spirv.LogicalEqual
 //===----------------------------------------------------------------------===//
 
 // CHECK-LABEL: @logical_equal_scalar
 spirv.func @logical_equal_scalar(%arg0: i1, %arg1: i1) "None" {
   // CHECK: llvm.icmp "eq" %{{.*}}, %{{.*}} : i1

mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls.mlir

-// RUN: mlir-opt --lower-host-to-llvm %s -split-input-file | FileCheck %s
+// RUN: mlir-opt --lower-host-to-llvm='use-opaque-pointers=1' %s -split-input-file | FileCheck %s
 
 module attributes {gpu.container_module, spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader], [SPV_KHR_variable_pointers]>, #spirv.resource_limits<max_compute_workgroup_invocations = 128, max_compute_workgroup_size = [128, 128, 64]>>} {
 
   //       CHECK: llvm.mlir.global linkonce @__spv__foo_bar_arg_0_descriptor_set0_binding0() {addr_space = 0 : i32} : !llvm.struct<(array<6 x i32>)>
   //       CHECK: llvm.func @__spv__foo_bar()
 
   //       CHECK: spirv.module @__spv__foo
   //       CHECK:   spirv.GlobalVariable @bar_arg_0 bind(0, 0) : !spirv.ptr<!spirv.struct<(!spirv.array<6 x i32, stride=4> [0])>, StorageBuffer>
   //       CHECK:   spirv.func @__spv__foo_bar
 
   //       CHECK:   spirv.EntryPoint "GLCompute" @__spv__foo_bar
   //       CHECK:   spirv.ExecutionMode @__spv__foo_bar "LocalSize", 1, 1, 1
 
   // CHECK-LABEL: @main
-  //       CHECK:   %[[SRC:.*]] = llvm.extractvalue %{{.*}}[0] : !llvm.struct<(ptr<i32>, ptr<i32>, i64, array<1 x i64>, array<1 x i64>)>
-  //  CHECK-NEXT:   %[[DEST:.*]] = llvm.mlir.addressof @__spv__foo_bar_arg_0_descriptor_set0_binding0 : !llvm.ptr<struct<(array<6 x i32>)>>
+  //       CHECK:   %[[SRC:.*]] = llvm.extractvalue %{{.*}}[0] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>
+  //  CHECK-NEXT:   %[[DEST:.*]] = llvm.mlir.addressof @__spv__foo_bar_arg_0_descriptor_set0_binding0 : !llvm.ptr
   //  CHECK-NEXT:   llvm.mlir.constant(false) : i1
-  //  CHECK-NEXT:   "llvm.intr.memcpy"(%[[DEST]], %[[SRC]], %[[SIZE:.*]], %{{.*}}) : (!llvm.ptr<struct<(array<6 x i32>)>>, !llvm.ptr<i32>, i64, i1) -> ()
+  //  CHECK-NEXT:   "llvm.intr.memcpy"(%[[DEST]], %[[SRC]], %[[SIZE:.*]], %{{.*}}) : (!llvm.ptr, !llvm.ptr, i64, i1) -> ()
   //  CHECK-NEXT:   llvm.call @__spv__foo_bar() : () -> ()
   //  CHECK-NEXT:   llvm.mlir.constant(false) : i1
-  //  CHECK-NEXT:   "llvm.intr.memcpy"(%[[SRC]], %[[DEST]], %[[SIZE]], %{{.*}}) : (!llvm.ptr<i32>, !llvm.ptr<struct<(array<6 x i32>)>>, i64, i1) -> ()
+  //  CHECK-NEXT:   "llvm.intr.memcpy"(%[[SRC]], %[[DEST]], %[[SIZE]], %{{.*}}) : (!llvm.ptr, !llvm.ptr, i64, i1) -> ()
 
   spirv.module @__spv__foo Logical GLSL450 requires #spirv.vce<v1.0, [Shader], [SPV_KHR_variable_pointers]> {
     spirv.GlobalVariable @bar_arg_0 bind(0, 0) : !spirv.ptr<!spirv.struct<(!spirv.array<6 x i32, stride=4> [0])>, StorageBuffer>
     spirv.func @bar() "None" attributes {workgroup_attributions = 0 : i64} {
       %0 = spirv.mlir.addressof @bar_arg_0 : !spirv.ptr<!spirv.struct<(!spirv.array<6 x i32, stride=4> [0])>, StorageBuffer>
       spirv.Return
     }
     spirv.EntryPoint "GLCompute" @bar

mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls_fail.mlir

-// RUN: mlir-opt --lower-host-to-llvm %s -verify-diagnostics
+// RUN: mlir-opt --lower-host-to-llvm='use-opaque-pointers=1' %s -verify-diagnostics
 
 module {
 // expected-error @+1 {{The module must contain exactly one entry point function}}
   spirv.module Logical GLSL450 {
   }
 }

mlir/test/Conversion/SPIRVToLLVM/memory-ops-to-llvm.mlir

-// RUN: mlir-opt -convert-spirv-to-llvm %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // spirv.AccessChain
 //===----------------------------------------------------------------------===//
 
 // CHECK-LABEL: @access_chain
 spirv.func @access_chain() "None" {
   // CHECK: %[[ONE:.*]] = llvm.mlir.constant(1 : i32) : i32
   %0 = spirv.Constant 1: i32
   %1 = spirv.Variable : !spirv.ptr<!spirv.struct<(f32, !spirv.array<4xf32>)>, Function>
   // CHECK: %[[ZERO:.*]] = llvm.mlir.constant(0 : i32) : i32
-  // CHECK: llvm.getelementptr %{{.*}}[%[[ZERO]], 1, %[[ONE]]] : (!llvm.ptr<struct<packed (f32, array<4 x f32>)>>, i32, i32) -> !llvm.ptr<f32>
+  // CHECK: llvm.getelementptr %{{.*}}[%[[ZERO]], 1, %[[ONE]]] : (!llvm.ptr, i32, i32) -> !llvm.ptr, !llvm.struct<packed (f32, array<4 x f32>)>
   %2 = spirv.AccessChain %1[%0, %0] : !spirv.ptr<!spirv.struct<(f32, !spirv.array<4xf32>)>, Function>, i32, i32
   spirv.Return
 }
 
 // CHECK-LABEL: @access_chain_array
 spirv.func @access_chain_array(%arg0 : i32) "None" {
   %0 = spirv.Variable : !spirv.ptr<!spirv.array<4x!spirv.array<4xf32>>, Function>
   // CHECK: %[[ZERO:.*]] = llvm.mlir.constant(0 : i32) : i32
-  // CHECK: llvm.getelementptr %{{.*}}[%[[ZERO]], %{{.*}}] : (!llvm.ptr<array<4 x array<4 x f32>>>, i32, i32) -> !llvm.ptr<array<4 x f32>>
+  // CHECK: llvm.getelementptr %{{.*}}[%[[ZERO]], %{{.*}}] : (!llvm.ptr, i32, i32) -> !llvm.ptr, !llvm.array<4 x array<4 x f32>>
   %1 = spirv.AccessChain %0[%arg0] : !spirv.ptr<!spirv.array<4x!spirv.array<4xf32>>, Function>, i32
   %2 = spirv.Load "Function" %1 ["Volatile"] : !spirv.array<4xf32>
   spirv.Return
 }
 
 //===----------------------------------------------------------------------===//
 // spirv.GlobalVariable and spirv.mlir.addressof
 //===----------------------------------------------------------------------===//
 
 spirv.module Logical GLSL450 {
   // CHECK: llvm.mlir.global external constant @var() {addr_space = 0 : i32} : f32
   spirv.GlobalVariable @var : !spirv.ptr<f32, Input>
 }
 
 spirv.module Logical GLSL450 {
   //       CHECK: llvm.mlir.global private @struct() {addr_space = 0 : i32} : !llvm.struct<packed (f32, array<10 x f32>)>
   // CHECK-LABEL: @func
-  //       CHECK:   llvm.mlir.addressof @struct : !llvm.ptr<struct<packed (f32, array<10 x f32>)>>
+  //       CHECK:   llvm.mlir.addressof @struct : !llvm.ptr
   spirv.GlobalVariable @struct : !spirv.ptr<!spirv.struct<(f32, !spirv.array<10xf32>)>, Private>
   spirv.func @func() "None" {
     %0 = spirv.mlir.addressof @struct : !spirv.ptr<!spirv.struct<(f32, !spirv.array<10xf32>)>, Private>
     spirv.Return
   }
 }
 
 spirv.module Logical GLSL450 {
   //       CHECK: llvm.mlir.global external @bar_descriptor_set0_binding0() {addr_space = 0 : i32} : i32
   // CHECK-LABEL: @foo
-  //       CHECK:   llvm.mlir.addressof @bar_descriptor_set0_binding0 : !llvm.ptr<i32>
+  //       CHECK:   llvm.mlir.addressof @bar_descriptor_set0_binding0 : !llvm.ptr
   spirv.GlobalVariable @bar bind(0, 0) : !spirv.ptr<i32, StorageBuffer>
   spirv.func @foo() "None" {
     %0 = spirv.mlir.addressof @bar : !spirv.ptr<i32, StorageBuffer>
     spirv.Return
   }
 }
 
 spirv.module @name Logical GLSL450 {
   //       CHECK: llvm.mlir.global external @name_bar_descriptor_set0_binding0() {addr_space = 0 : i32} : i32
   // CHECK-LABEL: @foo
-  //       CHECK:   llvm.mlir.addressof @name_bar_descriptor_set0_binding0 : !llvm.ptr<i32>
+  //       CHECK:   llvm.mlir.addressof @name_bar_descriptor_set0_binding0 : !llvm.ptr
   spirv.GlobalVariable @bar bind(0, 0) : !spirv.ptr<i32, StorageBuffer>
   spirv.func @foo() "None" {
     %0 = spirv.mlir.addressof @bar : !spirv.ptr<i32, StorageBuffer>
     spirv.Return
   }
 }
 
 spirv.module Logical GLSL450 {

 
 //===----------------------------------------------------------------------===//
 // spirv.Load
 //===----------------------------------------------------------------------===//
 
 // CHECK-LABEL: @load
 spirv.func @load() "None" {
   %0 = spirv.Variable : !spirv.ptr<f32, Function>
-  //  CHECK: llvm.load %{{.*}} : !llvm.ptr<f32>
+  //  CHECK: llvm.load %{{.*}} : !llvm.ptr -> f32
   %1 = spirv.Load "Function" %0 : f32
   spirv.Return
 }
 
 // CHECK-LABEL: @load_none
 spirv.func @load_none() "None" {
   %0 = spirv.Variable : !spirv.ptr<f32, Function>
-  //  CHECK: llvm.load %{{.*}} : !llvm.ptr<f32>
+  //  CHECK: llvm.load %{{.*}} : !llvm.ptr -> f32
   %1 = spirv.Load "Function" %0 ["None"] : f32
   spirv.Return
 }
 
 // CHECK-LABEL: @load_with_alignment
 spirv.func @load_with_alignment() "None" {
   %0 = spirv.Variable : !spirv.ptr<f32, Function>
-  // CHECK: llvm.load %{{.*}} {alignment = 4 : i64} : !llvm.ptr<f32>
+  // CHECK: llvm.load %{{.*}} {alignment = 4 : i64} : !llvm.ptr -> f32
   %1 = spirv.Load "Function" %0 ["Aligned", 4] : f32
   spirv.Return
 }
 
 // CHECK-LABEL: @load_volatile
 spirv.func @load_volatile() "None" {
   %0 = spirv.Variable : !spirv.ptr<f32, Function>
-  // CHECK: llvm.load volatile %{{.*}} : !llvm.ptr<f32>
+  // CHECK: llvm.load volatile %{{.*}} : !llvm.ptr -> f32
   %1 = spirv.Load "Function" %0 ["Volatile"] : f32
   spirv.Return
 }
 
 // CHECK-LABEL: @load_nontemporal
 spirv.func @load_nontemporal() "None" {
   %0 = spirv.Variable : !spirv.ptr<f32, Function>
-  // CHECK: llvm.load %{{.*}} {nontemporal} : !llvm.ptr<f32>
+  // CHECK: llvm.load %{{.*}} {nontemporal} : !llvm.ptr -> f32
   %1 = spirv.Load "Function" %0 ["Nontemporal"] : f32
   spirv.Return
 }
 
 //===----------------------------------------------------------------------===//
 // spirv.Store
 //===----------------------------------------------------------------------===//
 
 // CHECK-LABEL: @store
 spirv.func @store(%arg0 : f32) "None" {
   %0 = spirv.Variable : !spirv.ptr<f32, Function>
-  // CHECK: llvm.store %{{.*}}, %{{.*}} : !llvm.ptr<f32>
+  // CHECK: llvm.store %{{.*}}, %{{.*}} : f32, !llvm.ptr
   spirv.Store "Function" %0, %arg0 : f32
   spirv.Return
 }
 
 // CHECK-LABEL: @store_composite
 spirv.func @store_composite(%arg0 : !spirv.struct<(f64)>) "None" {
   %0 = spirv.Variable : !spirv.ptr<!spirv.struct<(f64)>, Function>
-  // CHECK: llvm.store %{{.*}}, %{{.*}} : !llvm.ptr<struct<packed (f64)>>
+  // CHECK: llvm.store %{{.*}}, %{{.*}} : !llvm.struct<packed (f64)>, !llvm.ptr
   spirv.Store "Function" %0, %arg0 : !spirv.struct<(f64)>
   spirv.Return
 }
 
 // CHECK-LABEL: @store_with_alignment
 spirv.func @store_with_alignment(%arg0 : f32) "None" {
   %0 = spirv.Variable : !spirv.ptr<f32, Function>
-  // CHECK: llvm.store %{{.*}}, %{{.*}} {alignment = 4 : i64} : !llvm.ptr<f32>
+  // CHECK: llvm.store %{{.*}}, %{{.*}} {alignment = 4 : i64} : f32, !llvm.ptr
   spirv.Store "Function" %0, %arg0 ["Aligned", 4] : f32
   spirv.Return
 }
 
 // CHECK-LABEL: @store_volatile
 spirv.func @store_volatile(%arg0 : f32) "None" {
   %0 = spirv.Variable : !spirv.ptr<f32, Function>
-  // CHECK: llvm.store volatile %{{.*}}, %{{.*}} : !llvm.ptr<f32>
+  // CHECK: llvm.store volatile %{{.*}}, %{{.*}} : f32, !llvm.ptr
   spirv.Store "Function" %0, %arg0 ["Volatile"] : f32
   spirv.Return
 }
 
 // CHECK-LABEL: @store_nontemporal
 spirv.func @store_nontemporal(%arg0 : f32) "None" {
   %0 = spirv.Variable : !spirv.ptr<f32, Function>
-  // CHECK: llvm.store %{{.*}}, %{{.*}} {nontemporal} : !llvm.ptr<f32>
+  // CHECK: llvm.store %{{.*}}, %{{.*}} {nontemporal} : f32, !llvm.ptr
   spirv.Store "Function" %0, %arg0 ["Nontemporal"] : f32
   spirv.Return
 }
 
 //===----------------------------------------------------------------------===//
 // spirv.Variable
 //===----------------------------------------------------------------------===//
 
 // CHECK-LABEL: @variable_scalar
 spirv.func @variable_scalar() "None" {
   // CHECK: %[[SIZE1:.*]] = llvm.mlir.constant(1 : i32) : i32
-  // CHECK: llvm.alloca %[[SIZE1]] x f32 : (i32) -> !llvm.ptr<f32>
+  // CHECK: llvm.alloca %[[SIZE1]] x f32 : (i32) -> !llvm.ptr
   %0 = spirv.Variable : !spirv.ptr<f32, Function>
   // CHECK: %[[SIZE2:.*]] = llvm.mlir.constant(1 : i32) : i32
-  // CHECK: llvm.alloca %[[SIZE2]] x i8 : (i32) -> !llvm.ptr<i8>
+  // CHECK: llvm.alloca %[[SIZE2]] x i8 : (i32) -> !llvm.ptr
   %1 = spirv.Variable : !spirv.ptr<i8, Function>
   spirv.Return
 }
 
 // CHECK-LABEL: @variable_scalar_with_initialization
 spirv.func @variable_scalar_with_initialization() "None" {
   // CHECK: %[[VALUE:.*]] = llvm.mlir.constant(0 : i64) : i64
   // CHECK: %[[SIZE:.*]] = llvm.mlir.constant(1 : i32) : i32
-  // CHECK: %[[ALLOCATED:.*]] = llvm.alloca %[[SIZE]] x i64 : (i32) -> !llvm.ptr<i64>
-  // CHECK: llvm.store %[[VALUE]], %[[ALLOCATED]] : !llvm.ptr<i64>
+  // CHECK: %[[ALLOCATED:.*]] = llvm.alloca %[[SIZE]] x i64 : (i32) -> !llvm.ptr
+  // CHECK: llvm.store %[[VALUE]], %[[ALLOCATED]] : i64, !llvm.ptr
   %c = spirv.Constant 0 : i64
   %0 = spirv.Variable init(%c) : !spirv.ptr<i64, Function>
   spirv.Return
 }
 
 // CHECK-LABEL: @variable_vector
 spirv.func @variable_vector() "None" {
   // CHECK: %[[SIZE:.*]] = llvm.mlir.constant(1 : i32) : i32
-  // CHECK: llvm.alloca  %[[SIZE]] x vector<3xf32> : (i32) -> !llvm.ptr<vector<3xf32>>
+  // CHECK: llvm.alloca  %[[SIZE]] x vector<3xf32> : (i32) -> !llvm.ptr
   %0 = spirv.Variable : !spirv.ptr<vector<3xf32>, Function>
   spirv.Return
 }
 
 // CHECK-LABEL: @variable_vector_with_initialization
 spirv.func @variable_vector_with_initialization() "None" {
   // CHECK: %[[VALUE:.*]] = llvm.mlir.constant(dense<false> : vector<3xi1>) : vector<3xi1>
   // CHECK: %[[SIZE:.*]] = llvm.mlir.constant(1 : i32) : i32
-  // CHECK: %[[ALLOCATED:.*]] = llvm.alloca %[[SIZE]] x vector<3xi1> : (i32) -> !llvm.ptr<vector<3xi1>>
-  // CHECK: llvm.store %[[VALUE]], %[[ALLOCATED]] : !llvm.ptr<vector<3xi1>>
+  // CHECK: %[[ALLOCATED:.*]] = llvm.alloca %[[SIZE]] x vector<3xi1> : (i32) -> !llvm.ptr
+  // CHECK: llvm.store %[[VALUE]], %[[ALLOCATED]] : vector<3xi1>, !llvm.ptr
   %c = spirv.Constant dense<false> : vector<3xi1>
   %0 = spirv.Variable init(%c) : !spirv.ptr<vector<3xi1>, Function>
   spirv.Return
 }
 
 // CHECK-LABEL: @variable_array
 spirv.func @variable_array() "None" {
   // CHECK: %[[SIZE:.*]] = llvm.mlir.constant(1 : i32) : i32
-  // CHECK: llvm.alloca %[[SIZE]] x !llvm.array<10 x i32> : (i32) -> !llvm.ptr<array<10 x i32>>
+  // CHECK: llvm.alloca %[[SIZE]] x !llvm.array<10 x i32> : (i32) -> !llvm.ptr
   %0 = spirv.Variable : !spirv.ptr<!spirv.array<10 x i32>, Function>
   spirv.Return
 }

mlir/test/Conversion/SPIRVToLLVM/misc-ops-to-llvm.mlir

-// RUN: mlir-opt -convert-spirv-to-llvm %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // spirv.CompositeExtract
 //===----------------------------------------------------------------------===//
 
 // CHECK-LABEL: @composite_extract_array
 spirv.func @composite_extract_array(%arg: !spirv.array<4x!spirv.array<4xf32>>) "None" {
   // CHECK: llvm.extractvalue %{{.*}}[1, 3] : !llvm.array<4 x array<4 x f32>>

mlir/test/Conversion/SPIRVToLLVM/module-ops-to-llvm.mlir

-// RUN: mlir-opt -convert-spirv-to-llvm %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // spirv.module
 //===----------------------------------------------------------------------===//
 
 // CHECK: module
 spirv.module Logical GLSL450 {}
 

mlir/test/Conversion/SPIRVToLLVM/shift-ops-to-llvm.mlir

-// RUN: mlir-opt -convert-spirv-to-llvm %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // spirv.ShiftRightArithmetic
 //===----------------------------------------------------------------------===//
 
 // CHECK-LABEL: @shift_right_arithmetic_scalar
 spirv.func @shift_right_arithmetic_scalar(%arg0: i32, %arg1: si32, %arg2 : i16, %arg3 : ui16) "None" {
   // CHECK: llvm.ashr %{{.*}}, %{{.*}} : i32

mlir/test/Conversion/SPIRVToLLVM/spirv-types-to-llvm-invalid.mlir

-// RUN: mlir-opt %s -convert-spirv-to-llvm -verify-diagnostics -split-input-file
+// RUN: mlir-opt %s -convert-spirv-to-llvm='use-opaque-pointers=1' -verify-diagnostics -split-input-file
 
 // expected-error@+1 {{failed to legalize operation 'spirv.func' that was explicitly marked illegal}}
 spirv.func @array_with_unnatural_stride(%arg: !spirv.array<4 x f32, stride=8>) -> () "None" {
   spirv.Return
 }
 
 // -----
 

mlir/test/Conversion/SPIRVToLLVM/spirv-types-to-llvm.mlir

-// RUN: mlir-opt -split-input-file -convert-spirv-to-llvm -verify-diagnostics %s | FileCheck %s
+// RUN: mlir-opt -split-input-file -convert-spirv-to-llvm='use-opaque-pointers=1' -verify-diagnostics %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // Array type
 //===----------------------------------------------------------------------===//
 
 // CHECK-LABEL: @array(!llvm.array<16 x f32>, !llvm.array<32 x vector<4xf32>>)
 spirv.func @array(!spirv.array<16 x f32>, !spirv.array< 32 x vector<4xf32> >) "None"
 
 // CHECK-LABEL: @array_with_natural_stride(!llvm.array<16 x f32>)
 spirv.func @array_with_natural_stride(!spirv.array<16 x f32, stride=4>) "None"
 
 //===----------------------------------------------------------------------===//
 // Pointer type
 //===----------------------------------------------------------------------===//
 
-// CHECK-LABEL: @pointer_scalar(!llvm.ptr<i1>, !llvm.ptr<f32>)
+// CHECK-LABEL: @pointer_scalar(!llvm.ptr, !llvm.ptr)
 spirv.func @pointer_scalar(!spirv.ptr<i1, Uniform>, !spirv.ptr<f32, Private>) "None"
 
-// CHECK-LABEL: @pointer_vector(!llvm.ptr<vector<4xi32>>)
+// CHECK-LABEL: @pointer_vector(!llvm.ptr)
 spirv.func @pointer_vector(!spirv.ptr<vector<4xi32>, Function>) "None"
 
 //===----------------------------------------------------------------------===//
 // Runtime array type
 //===----------------------------------------------------------------------===//
 
 // CHECK-LABEL: @runtime_array_vector(!llvm.array<0 x vector<4xf32>>)
 spirv.func @runtime_array_vector(!spirv.rtarray< vector<4xf32> >) "None"

mlir/test/Conversion/SPIRVToLLVM/typed-pointers.mlir

+// RUN: mlir-opt -split-input-file -convert-spirv-to-llvm='use-opaque-pointers=0' %s | FileCheck %s
+
+//===----------------------------------------------------------------------===//
+// Pointer type
+//===----------------------------------------------------------------------===//
+
+// CHECK-LABEL: @pointer_scalar(!llvm.ptr<i1>, !llvm.ptr<f32>)
+spirv.func @pointer_scalar(!spirv.ptr<i1, Uniform>, !spirv.ptr<f32, Private>) "None"
+
+// CHECK-LABEL: @pointer_vector(!llvm.ptr<vector<4xi32>>)
+spirv.func @pointer_vector(!spirv.ptr<vector<4xi32>, Function>) "None"
+
+// CHECK-LABEL: @bitcast_pointer
+spirv.func @bitcast_pointer(%arg0: !spirv.ptr<f32, Function>) "None" {
+  // CHECK: llvm.bitcast %{{.*}} : !llvm.ptr<f32> to !llvm.ptr<i32>
+  %0 = spirv.Bitcast %arg0 : !spirv.ptr<f32, Function> to !spirv.ptr<i32, Function>
+  spirv.Return
+}