Skip to content

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

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
Oct 30, 2023
Merged

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

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
3 changes: 0 additions & 3 deletions mlir/include/mlir/Conversion/Passes.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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",
Expand Down
4 changes: 2 additions & 2 deletions mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVM.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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();
}
Expand Down
1 change: 0 additions & 1 deletion mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -42,7 +42,6 @@ void ConvertSPIRVToLLVMPass::runOnOperation() {
ModuleOp module = getOperation();

LowerToLLVMOptions options(&getContext());
options.useOpaquePointers = useOpaquePointers;

LLVMTypeConverter converter(&getContext(), options);

Expand Down
2 changes: 1 addition & 1 deletion mlir/test/Conversion/SPIRVToLLVM/arithmetic-ops-to-llvm.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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
Expand Down
2 changes: 1 addition & 1 deletion mlir/test/Conversion/SPIRVToLLVM/bitwise-ops-to-llvm.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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
Expand Down
2 changes: 1 addition & 1 deletion mlir/test/Conversion/SPIRVToLLVM/cast-ops-to-llvm.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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
Expand Down
2 changes: 1 addition & 1 deletion mlir/test/Conversion/SPIRVToLLVM/comparison-ops-to-llvm.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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
Expand Down
2 changes: 1 addition & 1 deletion mlir/test/Conversion/SPIRVToLLVM/constant-op-to-llvm.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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
Expand Down
2 changes: 1 addition & 1 deletion mlir/test/Conversion/SPIRVToLLVM/control-flow-ops-to-llvm.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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
Expand Down
2 changes: 1 addition & 1 deletion mlir/test/Conversion/SPIRVToLLVM/func-ops-to-llvm.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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
Expand Down
2 changes: 1 addition & 1 deletion mlir/test/Conversion/SPIRVToLLVM/gl-ops-to-llvm.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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
Expand Down
2 changes: 1 addition & 1 deletion mlir/test/Conversion/SPIRVToLLVM/logical-ops-to-llvm.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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
Expand Down
2 changes: 1 addition & 1 deletion mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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]>>} {

Expand Down
2 changes: 1 addition & 1 deletion mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls_fail.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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}}
Expand Down
2 changes: 1 addition & 1 deletion mlir/test/Conversion/SPIRVToLLVM/memory-ops-to-llvm.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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
Expand Down
2 changes: 1 addition & 1 deletion mlir/test/Conversion/SPIRVToLLVM/misc-ops-to-llvm.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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
Expand Down
2 changes: 1 addition & 1 deletion mlir/test/Conversion/SPIRVToLLVM/module-ops-to-llvm.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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
Expand Down
2 changes: 1 addition & 1 deletion mlir/test/Conversion/SPIRVToLLVM/shift-ops-to-llvm.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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
Expand Down
4 changes: 2 additions & 2 deletions mlir/test/Conversion/SPIRVToLLVM/spirv-storage-class-mapping.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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)
Expand Down
2 changes: 1 addition & 1 deletion mlir/test/Conversion/SPIRVToLLVM/spirv-types-to-llvm-invalid.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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" {
Expand Down
2 changes: 1 addition & 1 deletion mlir/test/Conversion/SPIRVToLLVM/spirv-types-to-llvm.mlir
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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
Expand Down
18 changes: 0 additions & 18 deletions mlir/test/Conversion/SPIRVToLLVM/typed-pointers.mlir
View file
Open in desktop

This file was deleted.