[Mlir-commits] [mlir] 19b1e27 - [mlir][SPIRVToLLVM] Add pass option to emit opaque-pointers
Markus Böck
llvmlistbot at llvm.org
Mon Feb 13 13:23:54 PST 2023
Author: Markus Böck
Date: 2023-02-13T22:24:20+01:00
New Revision: 19b1e27fcd56a2a61525899a80bddffd8af9ba61
URL: https://github.com/llvm/llvm-project/commit/19b1e27fcd56a2a61525899a80bddffd8af9ba61
DIFF: https://github.com/llvm/llvm-project/commit/19b1e27fcd56a2a61525899a80bddffd8af9ba61.diff
LOG: [mlir][SPIRVToLLVM] Add pass option to emit opaque-pointers
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
Differential Revision: https://reviews.llvm.org/D143900
Added:
mlir/test/Conversion/SPIRVToLLVM/typed-pointers.mlir
Modified:
mlir/include/mlir/Conversion/Passes.td
mlir/lib/Conversion/SPIRVToLLVM/ConvertLaunchFuncToLLVMCalls.cpp
mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVM.cpp
mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.cpp
mlir/test/Conversion/SPIRVToLLVM/arithmetic-ops-to-llvm.mlir
mlir/test/Conversion/SPIRVToLLVM/bitwise-ops-to-llvm.mlir
mlir/test/Conversion/SPIRVToLLVM/cast-ops-to-llvm.mlir
mlir/test/Conversion/SPIRVToLLVM/comparison-ops-to-llvm.mlir
mlir/test/Conversion/SPIRVToLLVM/constant-op-to-llvm.mlir
mlir/test/Conversion/SPIRVToLLVM/control-flow-ops-to-llvm.mlir
mlir/test/Conversion/SPIRVToLLVM/func-ops-to-llvm.mlir
mlir/test/Conversion/SPIRVToLLVM/gl-ops-to-llvm.mlir
mlir/test/Conversion/SPIRVToLLVM/logical-ops-to-llvm.mlir
mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls.mlir
mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls_fail.mlir
mlir/test/Conversion/SPIRVToLLVM/memory-ops-to-llvm.mlir
mlir/test/Conversion/SPIRVToLLVM/misc-ops-to-llvm.mlir
mlir/test/Conversion/SPIRVToLLVM/module-ops-to-llvm.mlir
mlir/test/Conversion/SPIRVToLLVM/shift-ops-to-llvm.mlir
mlir/test/Conversion/SPIRVToLLVM/spirv-types-to-llvm-invalid.mlir
mlir/test/Conversion/SPIRVToLLVM/spirv-types-to-llvm.mlir
Removed:
################################################################################
diff --git a/mlir/include/mlir/Conversion/Passes.td b/mlir/include/mlir/Conversion/Passes.td
index 0533373b25e85..27fa50b6f971f 100644
--- a/mlir/include/mlir/Conversion/Passes.td
+++ b/mlir/include/mlir/Conversion/Passes.td
@@ -383,6 +383,12 @@ def LowerHostCodeToLLVMPass : Pass<"lower-host-to-llvm", "ModuleOp"> {
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"];
}
@@ -839,6 +845,12 @@ def ConvertSPIRVToLLVMPass : Pass<"convert-spirv-to-llvm", "ModuleOp"> {
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">
+ ];
}
//===----------------------------------------------------------------------===//
diff --git a/mlir/lib/Conversion/SPIRVToLLVM/ConvertLaunchFuncToLLVMCalls.cpp b/mlir/lib/Conversion/SPIRVToLLVM/ConvertLaunchFuncToLLVMCalls.cpp
index 7803593d73a74..08da805fe1295 100644
--- a/mlir/lib/Conversion/SPIRVToLLVM/ConvertLaunchFuncToLLVMCalls.cpp
+++ b/mlir/lib/Conversion/SPIRVToLLVM/ConvertLaunchFuncToLLVMCalls.cpp
@@ -262,7 +262,9 @@ class GPULaunchLowering : public ConvertOpToLLVMPattern<gpu::LaunchFuncOp> {
// 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;
@@ -302,6 +304,8 @@ class LowerHostCodeToLLVM
// 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);
diff --git a/mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVM.cpp b/mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVM.cpp
index 0ba294f2ed799..2cdce91806068 100644
--- a/mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVM.cpp
+++ b/mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVM.cpp
@@ -263,9 +263,9 @@ static std::optional<Type> convertArrayType(spirv::ArrayType type,
/// 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
@@ -317,8 +317,13 @@ class AccessChainPattern : public SPIRVToLLVMConversion<spirv::AccessChainOp> {
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();
}
};
@@ -1266,16 +1271,46 @@ class VariablePattern : public SPIRVToLLVMConversion<spirv::VariableOp> {
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
//===----------------------------------------------------------------------===//
@@ -1471,7 +1506,7 @@ void mlir::populateSPIRVToLLVMConversionPatterns(
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>,
diff --git a/mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.cpp b/mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.cpp
index 7766d8d9a0c9d..263276ef1b9b2 100644
--- a/mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.cpp
+++ b/mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.cpp
@@ -39,7 +39,11 @@ class ConvertSPIRVToLLVMPass
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);
diff --git a/mlir/test/Conversion/SPIRVToLLVM/arithmetic-ops-to-llvm.mlir b/mlir/test/Conversion/SPIRVToLLVM/arithmetic-ops-to-llvm.mlir
index dbbf8610afb4d..925443758c43d 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 %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' %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 a0afe0dafcaa2..af232bb2c6387 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 %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' %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 175274cf35344..4f6c1eaaf5048 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 %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' %s | FileCheck %s
//===----------------------------------------------------------------------===//
// spirv.Bitcast
@@ -41,7 +41,7 @@ spirv.func @bitcast_vector_to_vector(%arg0 : vector<4xf32>) "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>
+ // CHECK-NOT: llvm.bitcast
%0 = spirv.Bitcast %arg0 : !spirv.ptr<f32, Function> to !spirv.ptr<i32, Function>
spirv.Return
}
diff --git a/mlir/test/Conversion/SPIRVToLLVM/comparison-ops-to-llvm.mlir b/mlir/test/Conversion/SPIRVToLLVM/comparison-ops-to-llvm.mlir
index 52359db3be7bd..272ee309fdf6d 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 %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' %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 2d74022b34406..f66b0768a6143 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 %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' %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 3cb727e7dbaa4..8c58d59e86d7e 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 -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
diff --git a/mlir/test/Conversion/SPIRVToLLVM/func-ops-to-llvm.mlir b/mlir/test/Conversion/SPIRVToLLVM/func-ops-to-llvm.mlir
index 5b3d8ba5ca595..1dff2c48cefee 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 %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' %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 e1936e2fd8abe..851d164284c48 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 %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' %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 6d93480d3ed14..65db67314361f 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 %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' %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 550eca37cfb9d..348e3d5f85a3f 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 %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]>>} {
@@ -13,13 +13,13 @@ module attributes {gpu.container_module, spirv.target_env = #spirv.target_env<#s
// 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>
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 e36d30b434fc7..05ed608252613 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 %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}}
diff --git a/mlir/test/Conversion/SPIRVToLLVM/memory-ops-to-llvm.mlir b/mlir/test/Conversion/SPIRVToLLVM/memory-ops-to-llvm.mlir
index acdad2225ae16..04357a1b00a5b 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 %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' %s | FileCheck %s
//===----------------------------------------------------------------------===//
// spirv.AccessChain
@@ -10,7 +10,7 @@ spirv.func @access_chain() "None" {
%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
}
@@ -19,7 +19,7 @@ spirv.func @access_chain() "None" {
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
@@ -37,7 +37,7 @@ spirv.module Logical GLSL450 {
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>
@@ -48,7 +48,7 @@ spirv.module Logical GLSL450 {
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>
@@ -59,7 +59,7 @@ spirv.module Logical GLSL450 {
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>
@@ -94,7 +94,7 @@ spirv.module Logical GLSL450 {
// 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
}
@@ -102,7 +102,7 @@ spirv.func @load() "None" {
// 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
}
@@ -110,7 +110,7 @@ spirv.func @load_none() "None" {
// 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
}
@@ -118,7 +118,7 @@ spirv.func @load_with_alignment() "None" {
// 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
}
@@ -126,7 +126,7 @@ spirv.func @load_volatile() "None" {
// 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
}
@@ -138,7 +138,7 @@ spirv.func @load_nontemporal() "None" {
// 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
}
@@ -146,7 +146,7 @@ spirv.func @store(%arg0 : f32) "None" {
// 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
}
@@ -154,7 +154,7 @@ spirv.func @store_composite(%arg0 : !spirv.struct<(f64)>) "None" {
// 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
}
@@ -162,7 +162,7 @@ spirv.func @store_with_alignment(%arg0 : f32) "None" {
// 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
}
@@ -170,7 +170,7 @@ spirv.func @store_volatile(%arg0 : f32) "None" {
// 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
}
@@ -182,10 +182,10 @@ spirv.func @store_nontemporal(%arg0 : f32) "None" {
// 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
}
@@ -194,8 +194,8 @@ spirv.func @variable_scalar() "None" {
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
@@ -204,7 +204,7 @@ spirv.func @variable_scalar_with_initialization() "None" {
// 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
}
@@ -213,8 +213,8 @@ spirv.func @variable_vector() "None" {
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
@@ -223,7 +223,7 @@ spirv.func @variable_vector_with_initialization() "None" {
// 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
}
diff --git a/mlir/test/Conversion/SPIRVToLLVM/misc-ops-to-llvm.mlir b/mlir/test/Conversion/SPIRVToLLVM/misc-ops-to-llvm.mlir
index 13bde6e6fc563..b90022b25265d 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 %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' %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 894de8b9e90a9..d4cb88db720ca 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 %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' %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 da4def7aeda82..db8f207193c9a 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 %s | FileCheck %s
+// RUN: mlir-opt -convert-spirv-to-llvm='use-opaque-pointers=1' %s | FileCheck %s
//===----------------------------------------------------------------------===//
// spirv.ShiftRightArithmetic
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 084652aff9583..3965c47ec199f 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 -verify-diagnostics -split-input-file
+// RUN: mlir-opt %s -convert-spirv-to-llvm='use-opaque-pointers=1' -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 39038ad47f21a..167ad021a5fa1 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 -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
@@ -14,10 +14,10 @@ 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"
//===----------------------------------------------------------------------===//
diff --git a/mlir/test/Conversion/SPIRVToLLVM/typed-pointers.mlir b/mlir/test/Conversion/SPIRVToLLVM/typed-pointers.mlir
new file mode 100644
index 0000000000000..2c56f42a1fd52
--- /dev/null
+++ b/mlir/test/Conversion/SPIRVToLLVM/typed-pointers.mlir
@@ -0,0 +1,18 @@
+// 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