[MLIR][SPIRVToLLVM] Remove typed pointer support (#70568)

This commit removes the support for lowering SPIRV to LLVM dialect with
typed pointers. Typed pointers have been deprecated for a while now and
it's planned to soon remove them from the LLVM dialect.

Related PSA:
https://discourse.llvm.org/t/psa-removal-of-typed-pointers-from-the-llvm-dialect/74502

Author: Dinistro

mlir/include/mlir/Conversion/Passes.td

@@ -1016,9 +1016,6 @@ def ConvertSPIRVToLLVMPass : Pass<"convert-spirv-to-llvm", "ModuleOp"> {
   let dependentDialects = ["LLVM::LLVMDialect"];
 
   let options = [
-    Option<"useOpaquePointers", "use-opaque-pointers", "bool",
-                 /*default=*/"true", "Generate LLVM IR using opaque pointers "
-                 "instead of typed pointers">,
     Option<"clientAPI", "client-api", "::mlir::spirv::ClientAPI",
 	   /*default=*/"::mlir::spirv::ClientAPI::Unknown",
 	   "Derive StorageClass to address space mapping from the client API",

mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVM.cpp

@@ -1376,8 +1376,8 @@ class BitcastConversionPattern
     if (!dstType)
       return failure();
 
-    if (typeConverter.useOpaquePointers() &&
-        isa<LLVM::LLVMPointerType>(dstType)) {
+    // LLVM's opaque pointers do not require bitcasts.
+    if (isa<LLVM::LLVMPointerType>(dstType)) {
       rewriter.replaceOp(bitcastOp, adaptor.getOperand());
       return success();
     }

mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.cpp

@@ -42,7 +42,6 @@ void ConvertSPIRVToLLVMPass::runOnOperation() {
   ModuleOp module = getOperation();
 
   LowerToLLVMOptions options(&getContext());
-  options.useOpaquePointers = useOpaquePointers;
 
   LLVMTypeConverter converter(&getContext(), options);
 

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

@@ -1,4 +1,4 @@
-// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // spirv.IAdd

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

@@ -1,4 +1,4 @@
-// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // spirv.BitCount

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

@@ -1,4 +1,4 @@
-// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // spirv.Bitcast

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

@@ -1,4 +1,4 @@
-// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // spirv.IEqual

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

@@ -1,4 +1,4 @@
-// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // spirv.Constant

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

@@ -1,4 +1,4 @@
-// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' -split-input-file -verify-diagnostics %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm -split-input-file -verify-diagnostics %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // spirv.Branch

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

@@ -1,4 +1,4 @@
-// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // spirv.Return

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

@@ -1,4 +1,4 @@
-// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // spirv.GL.Ceil

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

@@ -1,4 +1,4 @@
-// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // spirv.LogicalEqual

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

@@ -1,4 +1,4 @@
-// RUN: mlir-opt --lower-host-to-llvm='use-opaque-pointers=1' %s -split-input-file | FileCheck %s
+// RUN: mlir-opt --lower-host-to-llvm %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]>>} {
 

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

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

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

@@ -1,4 +1,4 @@
-// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // spirv.AccessChain

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

@@ -1,4 +1,4 @@
-// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // spirv.CompositeExtract

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

@@ -1,4 +1,4 @@
-// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // spirv.module

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

@@ -1,4 +1,4 @@
-// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // spirv.ShiftRightArithmetic

mlir/test/Conversion/SPIRVToLLVM/spirv-storage-class-mapping.mlir

@@ -1,5 +1,5 @@
-// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' -verify-diagnostics %s | FileCheck %s --check-prefixes=CHECK-UNKNOWN,CHECK-ALL
-// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1 client-api=OpenCL' -verify-diagnostics %s | FileCheck %s --check-prefixes=CHECK-OPENCL,CHECK-ALL
+// RUN: mlir-opt -convert-spirv-to-llvm -verify-diagnostics %s | FileCheck %s --check-prefixes=CHECK-UNKNOWN,CHECK-ALL
+// RUN: mlir-opt -convert-spirv-to-llvm='client-api=OpenCL' -verify-diagnostics %s | FileCheck %s --check-prefixes=CHECK-OPENCL,CHECK-ALL
 
 // CHECK-OPENCL:         llvm.func @pointerUniformConstant(!llvm.ptr<2>)
 // CHECK-UNKNOWN:        llvm.func @pointerUniformConstant(!llvm.ptr)

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

@@ -1,4 +1,4 @@
-// RUN: mlir-opt %s -convert-spirv-to-llvm='use-opaque-pointers=1' -verify-diagnostics -split-input-file
+// RUN: mlir-opt %s -convert-spirv-to-llvm -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" {

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

@@ -1,4 +1,4 @@
-// RUN: mlir-opt -split-input-file -convert-spirv-to-llvm='use-opaque-pointers=1' -verify-diagnostics %s | FileCheck %s
+// RUN: mlir-opt -split-input-file -convert-spirv-to-llvm -verify-diagnostics %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // Array type