[Mlir-commits] [mlir] [MLIR][SPIRVToLLVM] Remove typed pointer support (PR #70568)
Christian Ulmann
llvmlistbot at llvm.org
Sat Oct 28 12:42:31 PDT 2023
https://github.com/Dinistro created https://github.com/llvm/llvm-project/pull/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.
I'm making this a draft for now, as we first require a PSA for this change.
>From d2415d637a4a9858b44b40f3c444dcf59d1be77f Mon Sep 17 00:00:00 2001
From: Christian Ulmann <christianulmann at gmail.com>
Date: Sat, 28 Oct 2023 21:36:41 +0200
Subject: [PATCH] [MLIR][SPIRVToLLVM] Remove typed pointer support
This commit removes the support for lowering SPIRV to LLVM dialect with
typed pointers. Typed pointers are deprecated for a while now and it's
planned to soon remove them from the LLVM dialect.
---
mlir/include/mlir/Conversion/Passes.td | 3 ---
.../lib/Conversion/SPIRVToLLVM/SPIRVToLLVM.cpp | 4 ++--
.../Conversion/SPIRVToLLVM/SPIRVToLLVMPass.cpp | 1 -
.../SPIRVToLLVM/arithmetic-ops-to-llvm.mlir | 2 +-
.../SPIRVToLLVM/bitwise-ops-to-llvm.mlir | 2 +-
.../SPIRVToLLVM/cast-ops-to-llvm.mlir | 2 +-
.../SPIRVToLLVM/comparison-ops-to-llvm.mlir | 2 +-
.../SPIRVToLLVM/constant-op-to-llvm.mlir | 2 +-
.../SPIRVToLLVM/control-flow-ops-to-llvm.mlir | 2 +-
.../SPIRVToLLVM/func-ops-to-llvm.mlir | 2 +-
.../Conversion/SPIRVToLLVM/gl-ops-to-llvm.mlir | 2 +-
.../SPIRVToLLVM/logical-ops-to-llvm.mlir | 2 +-
.../SPIRVToLLVM/lower-host-to-llvm-calls.mlir | 2 +-
.../lower-host-to-llvm-calls_fail.mlir | 2 +-
.../SPIRVToLLVM/memory-ops-to-llvm.mlir | 2 +-
.../SPIRVToLLVM/misc-ops-to-llvm.mlir | 2 +-
.../SPIRVToLLVM/module-ops-to-llvm.mlir | 2 +-
.../SPIRVToLLVM/shift-ops-to-llvm.mlir | 2 +-
.../spirv-storage-class-mapping.mlir | 4 ++--
.../spirv-types-to-llvm-invalid.mlir | 2 +-
.../SPIRVToLLVM/spirv-types-to-llvm.mlir | 2 +-
.../Conversion/SPIRVToLLVM/typed-pointers.mlir | 18 ------------------
22 files changed, 21 insertions(+), 43 deletions(-)
delete mode 100644 mlir/test/Conversion/SPIRVToLLVM/typed-pointers.mlir
diff --git a/mlir/include/mlir/Conversion/Passes.td b/mlir/include/mlir/Conversion/Passes.td
index 336f0d3af951b9a..cf6e545749ffc64 100644
--- a/mlir/include/mlir/Conversion/Passes.td
+++ b/mlir/include/mlir/Conversion/Passes.td
@@ -1019,9 +1019,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",
diff --git a/mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVM.cpp b/mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVM.cpp
index 60f34f413f587d4..a54163db9e1050e 100644
--- a/mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVM.cpp
+++ b/mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVM.cpp
@@ -1379,8 +1379,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();
}
diff --git a/mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.cpp b/mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.cpp
index 40798e9eb9dcbae..38091e449c56ee8 100644
--- a/mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.cpp
+++ b/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);
diff --git a/mlir/test/Conversion/SPIRVToLLVM/arithmetic-ops-to-llvm.mlir b/mlir/test/Conversion/SPIRVToLLVM/arithmetic-ops-to-llvm.mlir
index 925443758c43db1..dbbf8610afb4dfb 100644
--- a/mlir/test/Conversion/SPIRVToLLVM/arithmetic-ops-to-llvm.mlir
+++ b/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
diff --git a/mlir/test/Conversion/SPIRVToLLVM/bitwise-ops-to-llvm.mlir b/mlir/test/Conversion/SPIRVToLLVM/bitwise-ops-to-llvm.mlir
index af232bb2c6387d0..a0afe0dafcaa248 100644
--- a/mlir/test/Conversion/SPIRVToLLVM/bitwise-ops-to-llvm.mlir
+++ b/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
diff --git a/mlir/test/Conversion/SPIRVToLLVM/cast-ops-to-llvm.mlir b/mlir/test/Conversion/SPIRVToLLVM/cast-ops-to-llvm.mlir
index 4f6c1eaaf50487d..2026027a6dacee7 100644
--- a/mlir/test/Conversion/SPIRVToLLVM/cast-ops-to-llvm.mlir
+++ b/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
diff --git a/mlir/test/Conversion/SPIRVToLLVM/comparison-ops-to-llvm.mlir b/mlir/test/Conversion/SPIRVToLLVM/comparison-ops-to-llvm.mlir
index 272ee309fdf6d78..52359db3be7bdc0 100644
--- a/mlir/test/Conversion/SPIRVToLLVM/comparison-ops-to-llvm.mlir
+++ b/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
diff --git a/mlir/test/Conversion/SPIRVToLLVM/constant-op-to-llvm.mlir b/mlir/test/Conversion/SPIRVToLLVM/constant-op-to-llvm.mlir
index f66b0768a614360..2d74022b344066a 100644
--- a/mlir/test/Conversion/SPIRVToLLVM/constant-op-to-llvm.mlir
+++ b/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
diff --git a/mlir/test/Conversion/SPIRVToLLVM/control-flow-ops-to-llvm.mlir b/mlir/test/Conversion/SPIRVToLLVM/control-flow-ops-to-llvm.mlir
index 54ef71f75f528fc..3557830e779e240 100644
--- a/mlir/test/Conversion/SPIRVToLLVM/control-flow-ops-to-llvm.mlir
+++ b/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
diff --git a/mlir/test/Conversion/SPIRVToLLVM/func-ops-to-llvm.mlir b/mlir/test/Conversion/SPIRVToLLVM/func-ops-to-llvm.mlir
index 1dff2c48cefeead..5b3d8ba5ca59587 100644
--- a/mlir/test/Conversion/SPIRVToLLVM/func-ops-to-llvm.mlir
+++ b/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
diff --git a/mlir/test/Conversion/SPIRVToLLVM/gl-ops-to-llvm.mlir b/mlir/test/Conversion/SPIRVToLLVM/gl-ops-to-llvm.mlir
index 851d164284c48d9..e1936e2fd8abea0 100644
--- a/mlir/test/Conversion/SPIRVToLLVM/gl-ops-to-llvm.mlir
+++ b/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
diff --git a/mlir/test/Conversion/SPIRVToLLVM/logical-ops-to-llvm.mlir b/mlir/test/Conversion/SPIRVToLLVM/logical-ops-to-llvm.mlir
index 65db67314361f5a..6d93480d3ed142e 100644
--- a/mlir/test/Conversion/SPIRVToLLVM/logical-ops-to-llvm.mlir
+++ b/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
diff --git a/mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls.mlir b/mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls.mlir
index 71149f138940744..61944b9d047e49b 100644
--- a/mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls.mlir
+++ b/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]>>} {
diff --git a/mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls_fail.mlir b/mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls_fail.mlir
index 05ed60825261365..e36d30b434fc76e 100644
--- a/mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls_fail.mlir
+++ b/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}}
diff --git a/mlir/test/Conversion/SPIRVToLLVM/memory-ops-to-llvm.mlir b/mlir/test/Conversion/SPIRVToLLVM/memory-ops-to-llvm.mlir
index 04357a1b00a5b86..1847975b279afae 100644
--- a/mlir/test/Conversion/SPIRVToLLVM/memory-ops-to-llvm.mlir
+++ b/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
diff --git a/mlir/test/Conversion/SPIRVToLLVM/misc-ops-to-llvm.mlir b/mlir/test/Conversion/SPIRVToLLVM/misc-ops-to-llvm.mlir
index b90022b25265dc7..13bde6e6fc56305 100644
--- a/mlir/test/Conversion/SPIRVToLLVM/misc-ops-to-llvm.mlir
+++ b/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
diff --git a/mlir/test/Conversion/SPIRVToLLVM/module-ops-to-llvm.mlir b/mlir/test/Conversion/SPIRVToLLVM/module-ops-to-llvm.mlir
index d4cb88db720ca5f..894de8b9e90a924 100644
--- a/mlir/test/Conversion/SPIRVToLLVM/module-ops-to-llvm.mlir
+++ b/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
diff --git a/mlir/test/Conversion/SPIRVToLLVM/shift-ops-to-llvm.mlir b/mlir/test/Conversion/SPIRVToLLVM/shift-ops-to-llvm.mlir
index db8f207193c9a98..da4def7aeda82ae 100644
--- a/mlir/test/Conversion/SPIRVToLLVM/shift-ops-to-llvm.mlir
+++ b/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
diff --git a/mlir/test/Conversion/SPIRVToLLVM/spirv-storage-class-mapping.mlir b/mlir/test/Conversion/SPIRVToLLVM/spirv-storage-class-mapping.mlir
index 989ada93cf36ee7..b3991cbdbe8af1e 100644
--- a/mlir/test/Conversion/SPIRVToLLVM/spirv-storage-class-mapping.mlir
+++ b/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)
diff --git a/mlir/test/Conversion/SPIRVToLLVM/spirv-types-to-llvm-invalid.mlir b/mlir/test/Conversion/SPIRVToLLVM/spirv-types-to-llvm-invalid.mlir
index 3965c47ec199fcb..084652aff9583d6 100644
--- a/mlir/test/Conversion/SPIRVToLLVM/spirv-types-to-llvm-invalid.mlir
+++ b/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 at +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" {
diff --git a/mlir/test/Conversion/SPIRVToLLVM/spirv-types-to-llvm.mlir b/mlir/test/Conversion/SPIRVToLLVM/spirv-types-to-llvm.mlir
index 167ad021a5fa1dc..0f2dbf8ef115587 100644
--- a/mlir/test/Conversion/SPIRVToLLVM/spirv-types-to-llvm.mlir
+++ b/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
diff --git a/mlir/test/Conversion/SPIRVToLLVM/typed-pointers.mlir b/mlir/test/Conversion/SPIRVToLLVM/typed-pointers.mlir
deleted file mode 100644
index 2c56f42a1fd523a..000000000000000
--- a/mlir/test/Conversion/SPIRVToLLVM/typed-pointers.mlir
+++ /dev/null
@@ -1,18 +0,0 @@
-// 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
-}
More information about the Mlir-commits
mailing list