[Mlir-commits] [mlir] b0cd5b2 - [mlir][llvm] Switch remaining LLVM dialect tests to opaque pointers.
Tobias Gysi
llvmlistbot at llvm.org
Fri Mar 24 00:27:04 PDT 2023
Author: Tobias Gysi
Date: 2023-03-24T08:24:58+01:00
New Revision: b0cd5b2a476063b588c59325720c841d79ed3262
URL: https://github.com/llvm/llvm-project/commit/b0cd5b2a476063b588c59325720c841d79ed3262
DIFF: https://github.com/llvm/llvm-project/commit/b0cd5b2a476063b588c59325720c841d79ed3262.diff
LOG: [mlir][llvm] Switch remaining LLVM dialect tests to opaque pointers.
The revision switches the remaining LLVM dialect tests to use opaque
pointers. Selected tests are copied to a postfixed test file for the
time being.
A number of tests disappear once we fully switch to opaque pointers.
In particular, all tests that check verify a pointer element type
matches another type as well as tests of recursive types.
Part of https://discourse.llvm.org/t/rfc-switching-the-llvm-dialect-and-dialect-lowerings-to-opaque-pointers/68179
Reviewed By: Dinistro, zero9178
Differential Revision: https://reviews.llvm.org/D146726
Added:
mlir/test/Dialect/LLVMIR/canonicalize-typed-pointers.mlir
mlir/test/Dialect/LLVMIR/dynamic-gep-index-typed-pointers.mlir
mlir/test/Dialect/LLVMIR/global-typed-pointers.mlir
mlir/test/Dialect/LLVMIR/invalid-typed-pointers.mlir
mlir/test/Dialect/LLVMIR/layout-typed-pointers.mlir
mlir/test/Dialect/LLVMIR/nvvm-typed-pointers.mlir
mlir/test/Dialect/LLVMIR/parameter-attrs-invalid-typed-pointers.mlir
mlir/test/Dialect/LLVMIR/types-invalid-typed-pointers.mlir
mlir/test/Dialect/LLVMIR/types-typed-pointers.mlir
Modified:
mlir/test/Dialect/LLVMIR/callgraph.mlir
mlir/test/Dialect/LLVMIR/canonicalize.mlir
mlir/test/Dialect/LLVMIR/debuginfo.mlir
mlir/test/Dialect/LLVMIR/dynamic-gep-index.mlir
mlir/test/Dialect/LLVMIR/func.mlir
mlir/test/Dialect/LLVMIR/global.mlir
mlir/test/Dialect/LLVMIR/invalid.mlir
mlir/test/Dialect/LLVMIR/layout.mlir
mlir/test/Dialect/LLVMIR/nvvm.mlir
mlir/test/Dialect/LLVMIR/parameter-attrs-invalid.mlir
mlir/test/Dialect/LLVMIR/types-invalid.mlir
mlir/test/Dialect/LLVMIR/types.mlir
Removed:
################################################################################
diff --git a/mlir/test/Dialect/LLVMIR/callgraph.mlir b/mlir/test/Dialect/LLVMIR/callgraph.mlir
index edb5b35d126a..ca1044b8288c 100644
--- a/mlir/test/Dialect/LLVMIR/callgraph.mlir
+++ b/mlir/test/Dialect/LLVMIR/callgraph.mlir
@@ -58,33 +58,32 @@ module attributes {"test.name" = "Invoke call"} {
// CHECK-DAG: -- Call-Edge : <Unknown-Callee-Node>
// CHECK: -- SCCs --
- llvm.mlir.global external constant @_ZTIi() : !llvm.ptr<i8>
+ llvm.mlir.global external constant @_ZTIi() : !llvm.ptr
llvm.func @foo(%arg0: i32) -> !llvm.struct<(i32, f64, i32)>
- llvm.func @bar(!llvm.ptr<i8>, !llvm.ptr<i8>, !llvm.ptr<i8>)
+ llvm.func @bar(!llvm.ptr, !llvm.ptr, !llvm.ptr)
llvm.func @__gxx_personality_v0(...) -> i32
llvm.func @invokeLandingpad() -> i32 attributes { personality = @__gxx_personality_v0 } {
%0 = llvm.mlir.constant(0 : i32) : i32
%1 = llvm.mlir.constant(3 : i32) : i32
%2 = llvm.mlir.constant("\01") : !llvm.array<1 x i8>
- %3 = llvm.mlir.null : !llvm.ptr<ptr<i8>>
- %4 = llvm.mlir.null : !llvm.ptr<i8>
- %5 = llvm.mlir.addressof @_ZTIi : !llvm.ptr<ptr<i8>>
- %6 = llvm.bitcast %5 : !llvm.ptr<ptr<i8>> to !llvm.ptr<i8>
- %7 = llvm.mlir.constant(1 : i32) : i32
- %8 = llvm.alloca %7 x i8 : (i32) -> !llvm.ptr<i8>
- %9 = llvm.invoke @foo(%7) to ^bb2 unwind ^bb1 : (i32) -> !llvm.struct<(i32, f64, i32)>
+ %3 = llvm.mlir.null : !llvm.ptr
+ %4 = llvm.mlir.null : !llvm.ptr
+ %5 = llvm.mlir.addressof @_ZTIi : !llvm.ptr
+ %6 = llvm.mlir.constant(1 : i32) : i32
+ %7 = llvm.alloca %6 x i8 : (i32) -> !llvm.ptr
+ %8 = llvm.invoke @foo(%6) to ^bb2 unwind ^bb1 : (i32) -> !llvm.struct<(i32, f64, i32)>
^bb1:
- %10 = llvm.landingpad cleanup (catch %3 : !llvm.ptr<ptr<i8>>) (catch %6 : !llvm.ptr<i8>) (filter %2 : !llvm.array<1 x i8>) : !llvm.struct<(ptr<i8>, i32)>
- %11 = llvm.intr.eh.typeid.for %6 : (!llvm.ptr<i8>) -> i32
- llvm.resume %10 : !llvm.struct<(ptr<i8>, i32)>
+ %10 = llvm.landingpad cleanup (catch %3 : !llvm.ptr) (catch %5 : !llvm.ptr) (filter %2 : !llvm.array<1 x i8>) : !llvm.struct<(ptr, i32)>
+ %11 = llvm.intr.eh.typeid.for %5 : (!llvm.ptr) -> i32
+ llvm.resume %10 : !llvm.struct<(ptr, i32)>
^bb2:
- llvm.return %7 : i32
+ llvm.return %6 : i32
^bb3:
- llvm.invoke @bar(%8, %6, %4) to ^bb2 unwind ^bb1 : (!llvm.ptr<i8>, !llvm.ptr<i8>, !llvm.ptr<i8>) -> ()
+ llvm.invoke @bar(%7, %5, %4) to ^bb2 unwind ^bb1 : (!llvm.ptr, !llvm.ptr, !llvm.ptr) -> ()
^bb4:
llvm.return %0 : i32
diff --git a/mlir/test/Dialect/LLVMIR/canonicalize-typed-pointers.mlir b/mlir/test/Dialect/LLVMIR/canonicalize-typed-pointers.mlir
new file mode 100644
index 000000000000..2ae9727482fc
--- /dev/null
+++ b/mlir/test/Dialect/LLVMIR/canonicalize-typed-pointers.mlir
@@ -0,0 +1,86 @@
+// RUN: mlir-opt --pass-pipeline='builtin.module(llvm.func(canonicalize{test-convergence}))' %s -split-input-file | FileCheck %s
+
+// CHECK-LABEL: fold_bitcast
+// CHECK-SAME: %[[a0:arg[0-9]+]]
+// CHECK-NEXT: llvm.return %[[a0]]
+llvm.func @fold_bitcast(%x : !llvm.ptr<i8>) -> !llvm.ptr<i8> {
+ %c = llvm.bitcast %x : !llvm.ptr<i8> to !llvm.ptr<i8>
+ llvm.return %c : !llvm.ptr<i8>
+}
+
+// CHECK-LABEL: fold_bitcast2
+// CHECK-SAME: %[[a0:arg[0-9]+]]
+// CHECK-NEXT: llvm.return %[[a0]]
+llvm.func @fold_bitcast2(%x : !llvm.ptr<i8>) -> !llvm.ptr<i8> {
+ %c = llvm.bitcast %x : !llvm.ptr<i8> to !llvm.ptr<i32>
+ %d = llvm.bitcast %c : !llvm.ptr<i32> to !llvm.ptr<i8>
+ llvm.return %d : !llvm.ptr<i8>
+}
+
+// -----
+
+// CHECK-LABEL: fold_addrcast
+// CHECK-SAME: %[[a0:arg[0-9]+]]
+// CHECK-NEXT: llvm.return %[[a0]]
+llvm.func @fold_addrcast(%x : !llvm.ptr<i8>) -> !llvm.ptr<i8> {
+ %c = llvm.addrspacecast %x : !llvm.ptr<i8> to !llvm.ptr<i8>
+ llvm.return %c : !llvm.ptr<i8>
+}
+
+// CHECK-LABEL: fold_addrcast2
+// CHECK-SAME: %[[a0:arg[0-9]+]]
+// CHECK-NEXT: llvm.return %[[a0]]
+llvm.func @fold_addrcast2(%x : !llvm.ptr<i8>) -> !llvm.ptr<i8> {
+ %c = llvm.addrspacecast %x : !llvm.ptr<i8> to !llvm.ptr<i32, 5>
+ %d = llvm.addrspacecast %c : !llvm.ptr<i32, 5> to !llvm.ptr<i8>
+ llvm.return %d : !llvm.ptr<i8>
+}
+
+// -----
+
+// CHECK-LABEL: fold_gep
+// CHECK-SAME: %[[a0:arg[0-9]+]]
+// CHECK-NEXT: llvm.return %[[a0]]
+llvm.func @fold_gep(%x : !llvm.ptr<i8>) -> !llvm.ptr<i8> {
+ %c0 = arith.constant 0 : i32
+ %c = llvm.getelementptr %x[%c0] : (!llvm.ptr<i8>, i32) -> !llvm.ptr<i8>
+ llvm.return %c : !llvm.ptr<i8>
+}
+
+// -----
+
+// CHECK-LABEL: fold_gep_canon
+// CHECK-SAME: %[[a0:arg[0-9]+]]
+// CHECK-NEXT: %[[RES:.*]] = llvm.getelementptr %[[a0]][2]
+// CHECK-NEXT: llvm.return %[[RES]]
+llvm.func @fold_gep_canon(%x : !llvm.ptr<i8>) -> !llvm.ptr<i8> {
+ %c2 = arith.constant 2 : i32
+ %c = llvm.getelementptr %x[%c2] : (!llvm.ptr<i8>, i32) -> !llvm.ptr<i8>
+ llvm.return %c : !llvm.ptr<i8>
+}
+
+// -----
+
+// CHECK-LABEL: load_dce
+// CHECK-NEXT: llvm.return
+llvm.func @load_dce(%x : !llvm.ptr<i8>) {
+ %0 = llvm.load %x : !llvm.ptr<i8>
+ llvm.return
+}
+
+llvm.mlir.global external @fp() : !llvm.ptr<i8>
+
+// CHECK-LABEL: addr_dce
+// CHECK-NEXT: llvm.return
+llvm.func @addr_dce(%x : !llvm.ptr<i8>) {
+ %0 = llvm.mlir.addressof @fp : !llvm.ptr<ptr<i8>>
+ llvm.return
+}
+
+// CHECK-LABEL: alloca_dce
+// CHECK-NEXT: llvm.return
+llvm.func @alloca_dce() {
+ %c1_i64 = arith.constant 1 : i64
+ %0 = llvm.alloca %c1_i64 x i32 : (i64) -> !llvm.ptr<i32>
+ llvm.return
+}
diff --git a/mlir/test/Dialect/LLVMIR/canonicalize.mlir b/mlir/test/Dialect/LLVMIR/canonicalize.mlir
index 7fa7684f5ace..6b2cac14f298 100644
--- a/mlir/test/Dialect/LLVMIR/canonicalize.mlir
+++ b/mlir/test/Dialect/LLVMIR/canonicalize.mlir
@@ -37,8 +37,8 @@ llvm.func @no_fold_extractvalue(%arr: !llvm.array<4 x f32>) -> f32 {
%3 = llvm.extractvalue %2[0, 0] : !llvm.array<4 x !llvm.array<4 x f32>>
llvm.return %3 : f32
-
}
+
// -----
// CHECK-LABEL: fold_unrelated_extractvalue
@@ -56,18 +56,18 @@ llvm.func @fold_unrelated_extractvalue(%arr: !llvm.array<4 x f32>) -> f32 {
// CHECK-LABEL: fold_bitcast
// CHECK-SAME: %[[a0:arg[0-9]+]]
// CHECK-NEXT: llvm.return %[[a0]]
-llvm.func @fold_bitcast(%x : !llvm.ptr<i8>) -> !llvm.ptr<i8> {
- %c = llvm.bitcast %x : !llvm.ptr<i8> to !llvm.ptr<i8>
- llvm.return %c : !llvm.ptr<i8>
+llvm.func @fold_bitcast(%x : !llvm.ptr) -> !llvm.ptr {
+ %c = llvm.bitcast %x : !llvm.ptr to !llvm.ptr
+ llvm.return %c : !llvm.ptr
}
// CHECK-LABEL: fold_bitcast2
// CHECK-SAME: %[[a0:arg[0-9]+]]
// CHECK-NEXT: llvm.return %[[a0]]
-llvm.func @fold_bitcast2(%x : !llvm.ptr<i8>) -> !llvm.ptr<i8> {
- %c = llvm.bitcast %x : !llvm.ptr<i8> to !llvm.ptr<i32>
- %d = llvm.bitcast %c : !llvm.ptr<i32> to !llvm.ptr<i8>
- llvm.return %d : !llvm.ptr<i8>
+llvm.func @fold_bitcast2(%x : i32) -> i32 {
+ %c = llvm.bitcast %x : i32 to f32
+ %d = llvm.bitcast %c : f32 to i32
+ llvm.return %d : i32
}
// -----
@@ -75,18 +75,18 @@ llvm.func @fold_bitcast2(%x : !llvm.ptr<i8>) -> !llvm.ptr<i8> {
// CHECK-LABEL: fold_addrcast
// CHECK-SAME: %[[a0:arg[0-9]+]]
// CHECK-NEXT: llvm.return %[[a0]]
-llvm.func @fold_addrcast(%x : !llvm.ptr<i8>) -> !llvm.ptr<i8> {
- %c = llvm.addrspacecast %x : !llvm.ptr<i8> to !llvm.ptr<i8>
- llvm.return %c : !llvm.ptr<i8>
+llvm.func @fold_addrcast(%x : !llvm.ptr) -> !llvm.ptr {
+ %c = llvm.addrspacecast %x : !llvm.ptr to !llvm.ptr
+ llvm.return %c : !llvm.ptr
}
// CHECK-LABEL: fold_addrcast2
// CHECK-SAME: %[[a0:arg[0-9]+]]
// CHECK-NEXT: llvm.return %[[a0]]
-llvm.func @fold_addrcast2(%x : !llvm.ptr<i8>) -> !llvm.ptr<i8> {
- %c = llvm.addrspacecast %x : !llvm.ptr<i8> to !llvm.ptr<i32, 5>
- %d = llvm.addrspacecast %c : !llvm.ptr<i32, 5> to !llvm.ptr<i8>
- llvm.return %d : !llvm.ptr<i8>
+llvm.func @fold_addrcast2(%x : !llvm.ptr) -> !llvm.ptr {
+ %c = llvm.addrspacecast %x : !llvm.ptr to !llvm.ptr<5>
+ %d = llvm.addrspacecast %c : !llvm.ptr<5> to !llvm.ptr
+ llvm.return %d : !llvm.ptr
}
// -----
@@ -94,10 +94,10 @@ llvm.func @fold_addrcast2(%x : !llvm.ptr<i8>) -> !llvm.ptr<i8> {
// CHECK-LABEL: fold_gep
// CHECK-SAME: %[[a0:arg[0-9]+]]
// CHECK-NEXT: llvm.return %[[a0]]
-llvm.func @fold_gep(%x : !llvm.ptr<i8>) -> !llvm.ptr<i8> {
+llvm.func @fold_gep(%x : !llvm.ptr) -> !llvm.ptr {
%c0 = arith.constant 0 : i32
- %c = llvm.getelementptr %x[%c0] : (!llvm.ptr<i8>, i32) -> !llvm.ptr<i8>
- llvm.return %c : !llvm.ptr<i8>
+ %c = llvm.getelementptr %x[%c0] : (!llvm.ptr, i32) -> !llvm.ptr, i8
+ llvm.return %c : !llvm.ptr
}
// CHECK-LABEL: fold_gep_neg
@@ -114,13 +114,12 @@ llvm.func @fold_gep_neg(%x : !llvm.ptr) -> !llvm.ptr {
// CHECK-SAME: %[[a0:arg[0-9]+]]
// CHECK-NEXT: %[[RES:.*]] = llvm.getelementptr %[[a0]][2]
// CHECK-NEXT: llvm.return %[[RES]]
-llvm.func @fold_gep_canon(%x : !llvm.ptr<i8>) -> !llvm.ptr<i8> {
+llvm.func @fold_gep_canon(%x : !llvm.ptr) -> !llvm.ptr {
%c2 = arith.constant 2 : i32
- %c = llvm.getelementptr %x[%c2] : (!llvm.ptr<i8>, i32) -> !llvm.ptr<i8>
- llvm.return %c : !llvm.ptr<i8>
+ %c = llvm.getelementptr %x[%c2] : (!llvm.ptr, i32) -> !llvm.ptr, i8
+ llvm.return %c : !llvm.ptr
}
-
// -----
// Check that LLVM constants participate in cross-dialect constant folding. The
@@ -142,17 +141,17 @@ llvm.func @llvm_constant() -> i32 {
// CHECK-LABEL: load_dce
// CHECK-NEXT: llvm.return
-llvm.func @load_dce(%x : !llvm.ptr<i8>) {
- %0 = llvm.load %x : !llvm.ptr<i8>
+llvm.func @load_dce(%x : !llvm.ptr) {
+ %0 = llvm.load %x : !llvm.ptr -> i8
llvm.return
}
-llvm.mlir.global external @fp() : !llvm.ptr<i8>
+llvm.mlir.global external @fp() : !llvm.ptr
// CHECK-LABEL: addr_dce
// CHECK-NEXT: llvm.return
-llvm.func @addr_dce(%x : !llvm.ptr<i8>) {
- %0 = llvm.mlir.addressof @fp : !llvm.ptr<ptr<i8>>
+llvm.func @addr_dce(%x : !llvm.ptr) {
+ %0 = llvm.mlir.addressof @fp : !llvm.ptr
llvm.return
}
@@ -160,6 +159,6 @@ llvm.func @addr_dce(%x : !llvm.ptr<i8>) {
// CHECK-NEXT: llvm.return
llvm.func @alloca_dce() {
%c1_i64 = arith.constant 1 : i64
- %0 = llvm.alloca %c1_i64 x i32 : (i64) -> !llvm.ptr<i32>
+ %0 = llvm.alloca %c1_i64 x i32 : (i64) -> !llvm.ptr
llvm.return
}
diff --git a/mlir/test/Dialect/LLVMIR/debuginfo.mlir b/mlir/test/Dialect/LLVMIR/debuginfo.mlir
index 7aaef0d31bb9..f7517b2f2310 100644
--- a/mlir/test/Dialect/LLVMIR/debuginfo.mlir
+++ b/mlir/test/Dialect/LLVMIR/debuginfo.mlir
@@ -134,10 +134,10 @@
llvm.func @addr(%arg: i64) {
// CHECK: %[[ALLOC:.*]] = llvm.alloca
%allocCount = llvm.mlir.constant(1 : i32) : i32
- %alloc = llvm.alloca %allocCount x i64 : (i32) -> !llvm.ptr<i64>
+ %alloc = llvm.alloca %allocCount x i64 : (i32) -> !llvm.ptr
// CHECK: llvm.intr.dbg.declare #[[VAR0]] = %[[ALLOC]]
- llvm.intr.dbg.declare #var0 = %alloc : !llvm.ptr<i64>
+ llvm.intr.dbg.declare #var0 = %alloc : !llvm.ptr
llvm.return
}
diff --git a/mlir/test/Dialect/LLVMIR/dynamic-gep-index-typed-pointers.mlir b/mlir/test/Dialect/LLVMIR/dynamic-gep-index-typed-pointers.mlir
new file mode 100644
index 000000000000..9e14b1db3432
--- /dev/null
+++ b/mlir/test/Dialect/LLVMIR/dynamic-gep-index-typed-pointers.mlir
@@ -0,0 +1,12 @@
+// RUN: mlir-opt %s | FileCheck %s
+
+module attributes {dlti.dl_spec = #dlti.dl_spec<#dlti.dl_entry<i1, dense<8> : vector<2xi32>>, #dlti.dl_entry<i8, dense<8> : vector<2xi32>>, #dlti.dl_entry<i16, dense<16> : vector<2xi32>>, #dlti.dl_entry<i32, dense<32> : vector<2xi32>>, #dlti.dl_entry<i64, dense<[32, 64]> : vector<2xi32>>, #dlti.dl_entry<f16, dense<16> : vector<2xi32>>, #dlti.dl_entry<f64, dense<64> : vector<2xi32>>, #dlti.dl_entry<f128, dense<128> : vector<2xi32>>>} {
+ // CHECK: llvm.func @foo(%[[ARG0:.+]]: !llvm.ptr<struct<"my_struct", {{.+}}>>, %[[ARG1:.+]]: i32)
+ llvm.func @foo(%arg0: !llvm.ptr<struct<"my_struct", (struct<"sub_struct", (i32, i8)>, array<4 x i32>)>>, %arg1: i32) {
+ // CHECK: %[[C0:.+]] = llvm.mlir.constant(0 : i32)
+ %0 = llvm.mlir.constant(0 : i32) : i32
+ // CHECK: llvm.getelementptr %[[ARG0]][%[[C0]], 1, %[[ARG1]]]
+ %1 = "llvm.getelementptr"(%arg0, %0, %arg1) {rawConstantIndices = array<i32: -2147483648, 1, -2147483648>} : (!llvm.ptr<struct<"my_struct", (struct<"sub_struct", (i32, i8)>, array<4 x i32>)>>, i32, i32) -> !llvm.ptr<i32>
+ llvm.return
+ }
+}
diff --git a/mlir/test/Dialect/LLVMIR/dynamic-gep-index.mlir b/mlir/test/Dialect/LLVMIR/dynamic-gep-index.mlir
index 9e14b1db3432..f5808134ea02 100644
--- a/mlir/test/Dialect/LLVMIR/dynamic-gep-index.mlir
+++ b/mlir/test/Dialect/LLVMIR/dynamic-gep-index.mlir
@@ -1,12 +1,12 @@
// RUN: mlir-opt %s | FileCheck %s
module attributes {dlti.dl_spec = #dlti.dl_spec<#dlti.dl_entry<i1, dense<8> : vector<2xi32>>, #dlti.dl_entry<i8, dense<8> : vector<2xi32>>, #dlti.dl_entry<i16, dense<16> : vector<2xi32>>, #dlti.dl_entry<i32, dense<32> : vector<2xi32>>, #dlti.dl_entry<i64, dense<[32, 64]> : vector<2xi32>>, #dlti.dl_entry<f16, dense<16> : vector<2xi32>>, #dlti.dl_entry<f64, dense<64> : vector<2xi32>>, #dlti.dl_entry<f128, dense<128> : vector<2xi32>>>} {
- // CHECK: llvm.func @foo(%[[ARG0:.+]]: !llvm.ptr<struct<"my_struct", {{.+}}>>, %[[ARG1:.+]]: i32)
- llvm.func @foo(%arg0: !llvm.ptr<struct<"my_struct", (struct<"sub_struct", (i32, i8)>, array<4 x i32>)>>, %arg1: i32) {
+ // CHECK: llvm.func @foo(%[[ARG0:.+]]: !llvm.ptr, %[[ARG1:.+]]: i32)
+ llvm.func @foo(%arg0: !llvm.ptr, %arg1: i32) {
// CHECK: %[[C0:.+]] = llvm.mlir.constant(0 : i32)
%0 = llvm.mlir.constant(0 : i32) : i32
// CHECK: llvm.getelementptr %[[ARG0]][%[[C0]], 1, %[[ARG1]]]
- %1 = "llvm.getelementptr"(%arg0, %0, %arg1) {rawConstantIndices = array<i32: -2147483648, 1, -2147483648>} : (!llvm.ptr<struct<"my_struct", (struct<"sub_struct", (i32, i8)>, array<4 x i32>)>>, i32, i32) -> !llvm.ptr<i32>
+ %1 = "llvm.getelementptr"(%arg0, %0, %arg1) {elem_type = !llvm.struct<"my_struct", (struct<"sub_struct", (i32, i8)>, array<4 x i32>)>, rawConstantIndices = array<i32: -2147483648, 1, -2147483648>} : (!llvm.ptr, i32, i32) -> !llvm.ptr
llvm.return
}
}
diff --git a/mlir/test/Dialect/LLVMIR/func.mlir b/mlir/test/Dialect/LLVMIR/func.mlir
index 5cc7d75b627f..50f6c6a0e56f 100644
--- a/mlir/test/Dialect/LLVMIR/func.mlir
+++ b/mlir/test/Dialect/LLVMIR/func.mlir
@@ -33,10 +33,10 @@ module {
// GENERIC-SAME: () -> ()
}) {sym_name = "baz", function_type = !llvm.func<i64 (i64)>} : () -> ()
- // CHECK: llvm.func @qux(!llvm.ptr<i64> {llvm.noalias}, i64)
+ // CHECK: llvm.func @qux(!llvm.ptr {llvm.noalias}, i64)
// CHECK: attributes {xxx = {yyy = 42 : i64}}
"llvm.func"() ({
- }) {sym_name = "qux", function_type = !llvm.func<void (ptr<i64>, i64)>,
+ }) {sym_name = "qux", function_type = !llvm.func<void (ptr, i64)>,
arg_attrs = [{llvm.noalias}, {}], xxx = {yyy = 42}} : () -> ()
// CHECK: llvm.func @roundtrip1()
@@ -71,56 +71,56 @@ module {
// CHECK: llvm.func @roundtrip8() -> i32
llvm.func @roundtrip8() -> i32 attributes {}
- // CHECK: llvm.func @roundtrip9(!llvm.ptr<i32> {llvm.noalias})
- llvm.func @roundtrip9(!llvm.ptr<i32> {llvm.noalias})
+ // CHECK: llvm.func @roundtrip9(!llvm.ptr {llvm.noalias})
+ llvm.func @roundtrip9(!llvm.ptr {llvm.noalias})
- // CHECK: llvm.func @roundtrip10(!llvm.ptr<i32> {llvm.noalias})
- llvm.func @roundtrip10(%arg0: !llvm.ptr<i32> {llvm.noalias})
+ // CHECK: llvm.func @roundtrip10(!llvm.ptr {llvm.noalias})
+ llvm.func @roundtrip10(%arg0: !llvm.ptr {llvm.noalias})
- // CHECK: llvm.func @roundtrip11(%{{.*}}: !llvm.ptr<i32> {llvm.noalias}) {
- llvm.func @roundtrip11(%arg0: !llvm.ptr<i32> {llvm.noalias}) {
+ // CHECK: llvm.func @roundtrip11(%{{.*}}: !llvm.ptr {llvm.noalias}) {
+ llvm.func @roundtrip11(%arg0: !llvm.ptr {llvm.noalias}) {
llvm.return
}
- // CHECK: llvm.func @roundtrip12(%{{.*}}: !llvm.ptr<i32> {llvm.noalias})
+ // CHECK: llvm.func @roundtrip12(%{{.*}}: !llvm.ptr {llvm.noalias})
// CHECK: attributes {foo = 42 : i32}
- llvm.func @roundtrip12(%arg0: !llvm.ptr<i32> {llvm.noalias})
+ llvm.func @roundtrip12(%arg0: !llvm.ptr {llvm.noalias})
attributes {foo = 42 : i32} {
llvm.return
}
- // CHECK: llvm.func @byvalattr(%{{.*}}: !llvm.ptr<i32> {llvm.byval = i32})
- llvm.func @byvalattr(%arg0: !llvm.ptr<i32> {llvm.byval = i32}) {
+ // CHECK: llvm.func @byvalattr(%{{.*}}: !llvm.ptr {llvm.byval = i32})
+ llvm.func @byvalattr(%arg0: !llvm.ptr {llvm.byval = i32}) {
llvm.return
}
- // CHECK: llvm.func @sretattr(%{{.*}}: !llvm.ptr<i32> {llvm.sret = i32})
- // LOCINFO: llvm.func @sretattr(%{{.*}}: !llvm.ptr<i32> {llvm.sret = i32} loc("some_source_loc"))
- llvm.func @sretattr(%arg0: !llvm.ptr<i32> {llvm.sret = i32} loc("some_source_loc")) {
+ // CHECK: llvm.func @sretattr(%{{.*}}: !llvm.ptr {llvm.sret = i32})
+ // LOCINFO: llvm.func @sretattr(%{{.*}}: !llvm.ptr {llvm.sret = i32} loc("some_source_loc"))
+ llvm.func @sretattr(%arg0: !llvm.ptr {llvm.sret = i32} loc("some_source_loc")) {
llvm.return
}
- // CHECK: llvm.func @nestattr(%{{.*}}: !llvm.ptr<i32> {llvm.nest})
- llvm.func @nestattr(%arg0: !llvm.ptr<i32> {llvm.nest}) {
+ // CHECK: llvm.func @nestattr(%{{.*}}: !llvm.ptr {llvm.nest})
+ llvm.func @nestattr(%arg0: !llvm.ptr {llvm.nest}) {
llvm.return
}
- // CHECK: llvm.func @llvm_noalias_decl(!llvm.ptr<f32> {llvm.noalias})
- llvm.func @llvm_noalias_decl(!llvm.ptr<f32> {llvm.noalias})
- // CHECK: llvm.func @byrefattr_decl(!llvm.ptr<i32> {llvm.byref = i32})
- llvm.func @byrefattr_decl(!llvm.ptr<i32> {llvm.byref = i32})
- // CHECK: llvm.func @byvalattr_decl(!llvm.ptr<i32> {llvm.byval = i32})
- llvm.func @byvalattr_decl(!llvm.ptr<i32> {llvm.byval = i32})
- // CHECK: llvm.func @sretattr_decl(!llvm.ptr<i32> {llvm.sret = i32})
- llvm.func @sretattr_decl(!llvm.ptr<i32> {llvm.sret = i32})
- // CHECK: llvm.func @nestattr_decl(!llvm.ptr<i32> {llvm.nest})
- llvm.func @nestattr_decl(!llvm.ptr<i32> {llvm.nest})
+ // CHECK: llvm.func @llvm_noalias_decl(!llvm.ptr {llvm.noalias})
+ llvm.func @llvm_noalias_decl(!llvm.ptr {llvm.noalias})
+ // CHECK: llvm.func @byrefattr_decl(!llvm.ptr {llvm.byref = i32})
+ llvm.func @byrefattr_decl(!llvm.ptr {llvm.byref = i32})
+ // CHECK: llvm.func @byvalattr_decl(!llvm.ptr {llvm.byval = i32})
+ llvm.func @byvalattr_decl(!llvm.ptr {llvm.byval = i32})
+ // CHECK: llvm.func @sretattr_decl(!llvm.ptr {llvm.sret = i32})
+ llvm.func @sretattr_decl(!llvm.ptr {llvm.sret = i32})
+ // CHECK: llvm.func @nestattr_decl(!llvm.ptr {llvm.nest})
+ llvm.func @nestattr_decl(!llvm.ptr {llvm.nest})
// CHECK: llvm.func @noundefattr_decl(i32 {llvm.noundef})
llvm.func @noundefattr_decl(i32 {llvm.noundef})
- // CHECK: llvm.func @llvm_align_decl(!llvm.ptr<f32> {llvm.align = 4 : i64})
- llvm.func @llvm_align_decl(!llvm.ptr<f32> {llvm.align = 4})
- // CHECK: llvm.func @inallocaattr_decl(!llvm.ptr<i32> {llvm.inalloca = i32})
- llvm.func @inallocaattr_decl(!llvm.ptr<i32> {llvm.inalloca = i32})
+ // CHECK: llvm.func @llvm_align_decl(!llvm.ptr {llvm.align = 4 : i64})
+ llvm.func @llvm_align_decl(!llvm.ptr {llvm.align = 4})
+ // CHECK: llvm.func @inallocaattr_decl(!llvm.ptr {llvm.inalloca = i32})
+ llvm.func @inallocaattr_decl(!llvm.ptr {llvm.inalloca = i32})
// CHECK: llvm.func @variadic(...)
diff --git a/mlir/test/Dialect/LLVMIR/global-typed-pointers.mlir b/mlir/test/Dialect/LLVMIR/global-typed-pointers.mlir
new file mode 100644
index 000000000000..56d720cc866b
--- /dev/null
+++ b/mlir/test/Dialect/LLVMIR/global-typed-pointers.mlir
@@ -0,0 +1,46 @@
+// RUN: mlir-opt -split-input-file -verify-diagnostics %s | FileCheck %s
+
+// CHECK: llvm.mlir.global internal @global(42 : i64) {addr_space = 0 : i32} : i64
+llvm.mlir.global internal @global(42 : i64) : i64
+
+// CHECK: llvm.mlir.global internal constant @".string"("foobar")
+llvm.mlir.global internal constant @".string"("foobar") : !llvm.array<6 x i8>
+
+func.func @references() {
+ // CHECK: llvm.mlir.addressof @global : !llvm.ptr<i64>
+ %0 = llvm.mlir.addressof @global : !llvm.ptr<i64>
+
+ // CHECK: llvm.mlir.addressof @".string" : !llvm.ptr<array<6 x i8>>
+ %1 = llvm.mlir.addressof @".string" : !llvm.ptr<array<6 x i8>>
+
+ llvm.return
+}
+
+// -----
+
+llvm.mlir.global internal @foo(0: i32) : i32
+
+func.func @bar() {
+ // expected-error @+1 {{the type must be a pointer to the type of the referenced global}}
+ llvm.mlir.addressof @foo : !llvm.ptr<i64>
+ llvm.return
+}
+
+// -----
+
+llvm.func @foo()
+
+llvm.func @bar() {
+ // expected-error @+1 {{the type must be a pointer to the type of the referenced function}}
+ llvm.mlir.addressof @foo : !llvm.ptr<i8>
+ llvm.return
+}
+
+// -----
+
+llvm.mlir.global internal @g(32 : i64) {addr_space = 3: i32} : i64
+func.func @mismatch_addr_space() {
+ // expected-error @+1 {{pointer address space must match address space of the referenced global}}
+ llvm.mlir.addressof @g : !llvm.ptr<i64, 4>
+ llvm.return
+}
diff --git a/mlir/test/Dialect/LLVMIR/global.mlir b/mlir/test/Dialect/LLVMIR/global.mlir
index 2f0850834a0e..aff116db5dcc 100644
--- a/mlir/test/Dialect/LLVMIR/global.mlir
+++ b/mlir/test/Dialect/LLVMIR/global.mlir
@@ -66,17 +66,14 @@ llvm.mlir.global external @has_addr_space(32 : i64) {addr_space = 3: i32} : i64
// CHECK-LABEL: references
func.func @references() {
- // CHECK: llvm.mlir.addressof @global : !llvm.ptr<i64>
- %0 = llvm.mlir.addressof @global : !llvm.ptr<i64>
-
- // CHECK: llvm.mlir.addressof @".string" : !llvm.ptr<array<6 x i8>>
- %1 = llvm.mlir.addressof @".string" : !llvm.ptr<array<6 x i8>>
+ // CHECK: llvm.mlir.addressof @".string" : !llvm.ptr
+ %0 = llvm.mlir.addressof @".string" : !llvm.ptr
// CHECK: llvm.mlir.addressof @global : !llvm.ptr
- %2 = llvm.mlir.addressof @global : !llvm.ptr
+ %1 = llvm.mlir.addressof @global : !llvm.ptr
// CHECK: llvm.mlir.addressof @has_addr_space : !llvm.ptr<3>
- %3 = llvm.mlir.addressof @has_addr_space : !llvm.ptr<3>
+ %2 = llvm.mlir.addressof @has_addr_space : !llvm.ptr<3>
llvm.return
}
@@ -164,7 +161,7 @@ func.func @foo() {
// The attribute parser will consume the first colon-type, so we put two of
// them to trigger the attribute type mismatch error.
// expected-error @+1 {{invalid kind of attribute specified}}
- llvm.mlir.addressof "foo" : i64 : !llvm.ptr<func<void ()>>
+ llvm.mlir.addressof "foo" : i64 : !llvm.ptr
llvm.return
}
@@ -172,27 +169,7 @@ func.func @foo() {
func.func @foo() {
// expected-error @+1 {{must reference a global defined by 'llvm.mlir.global'}}
- llvm.mlir.addressof @foo : !llvm.ptr<func<void ()>>
- llvm.return
-}
-
-// -----
-
-llvm.mlir.global internal @foo(0: i32) : i32
-
-func.func @bar() {
- // expected-error @+1 {{the type must be a pointer to the type of the referenced global}}
- llvm.mlir.addressof @foo : !llvm.ptr<i64>
- llvm.return
-}
-
-// -----
-
-llvm.func @foo()
-
-llvm.func @bar() {
- // expected-error @+1 {{the type must be a pointer to the type of the referenced function}}
- llvm.mlir.addressof @foo : !llvm.ptr<i8>
+ llvm.mlir.addressof @foo : !llvm.ptr
llvm.return
}
@@ -224,23 +201,15 @@ llvm.mlir.global internal @g(43 : i64) : i64 {
llvm.mlir.global internal @g(32 : i64) {addr_space = 3: i32} : i64
func.func @mismatch_addr_space_implicit_global() {
// expected-error @+1 {{pointer address space must match address space of the referenced global}}
- llvm.mlir.addressof @g : !llvm.ptr<i64>
+ llvm.mlir.addressof @g : !llvm.ptr
llvm.return
}
// -----
llvm.mlir.global internal @g(32 : i64) {addr_space = 3: i32} : i64
-func.func @mismatch_addr_space() {
- // expected-error @+1 {{pointer address space must match address space of the referenced global}}
- llvm.mlir.addressof @g : !llvm.ptr<i64, 4>
- llvm.return
-}
-// -----
-llvm.mlir.global internal @g(32 : i64) {addr_space = 3: i32} : i64
-
-func.func @mismatch_addr_space_opaque() {
+func.func @mismatch_addr_space() {
// expected-error @+1 {{pointer address space must match address space of the referenced global}}
llvm.mlir.addressof @g : !llvm.ptr<4>
llvm.return
diff --git a/mlir/test/Dialect/LLVMIR/invalid-typed-pointers.mlir b/mlir/test/Dialect/LLVMIR/invalid-typed-pointers.mlir
new file mode 100644
index 000000000000..033b84d04ef8
--- /dev/null
+++ b/mlir/test/Dialect/LLVMIR/invalid-typed-pointers.mlir
@@ -0,0 +1,283 @@
+// RUN: mlir-opt -allow-unregistered-dialect %s -split-input-file -verify-diagnostics
+
+func.func @alloca_ptr_type_attr_non_opaque_ptr(%sz : i64) {
+ // expected-error at below {{unexpected 'elem_type' attribute when non-opaque pointer type is used}}
+ "llvm.alloca"(%sz) { elem_type = i32 } : (i64) -> !llvm.ptr<i32>
+}
+
+// -----
+
+func.func @gep_missing_input_type(%pos : i64, %base : !llvm.ptr<f32>) {
+ // expected-error at +1 {{2 operands present, but expected 0}}
+ llvm.getelementptr %base[%pos] : () -> (!llvm.ptr<f32>)
+}
+
+// -----
+
+func.func @gep_missing_result_type(%pos : i64, %base : !llvm.ptr<f32>) {
+ // expected-error at +1 {{op requires one result}}
+ llvm.getelementptr %base[%pos] : (!llvm.ptr<f32>, i64) -> ()
+}
+
+// -----
+
+func.func @gep_non_function_type(%pos : i64, %base : !llvm.ptr<f32>) {
+ // expected-error at +1 {{invalid kind of type specified}}
+ llvm.getelementptr %base[%pos] : !llvm.ptr<f32>
+}
+
+// -----
+
+func.func @gep_too_few_dynamic(%base : !llvm.ptr<f32>) {
+ // expected-error at +1 {{expected as many dynamic indices as specified in 'rawConstantIndices'}}
+ %1 = "llvm.getelementptr"(%base) {rawConstantIndices = array<i32: -2147483648>} : (!llvm.ptr<f32>) -> !llvm.ptr<f32>
+}
+
+// -----
+
+func.func @call_variadic(%callee : !llvm.ptr<func<i8 (i8, ...)>>, %arg : i8) {
+ // expected-error at +1 {{indirect calls to variadic functions are not supported}}
+ llvm.call %callee(%arg) : !llvm.ptr<func<i8 (i8, ...)>>, (i8) -> (i8)
+ llvm.return
+}
+
+// -----
+
+func.func @indirect_callee_arg_mismatch(%arg0 : i32, %callee : !llvm.ptr<func<void(i8)>>) {
+ // expected-error at +1 {{'llvm.call' op operand type mismatch for operand 0: 'i32' != 'i8'}}
+ "llvm.call"(%callee, %arg0) : (!llvm.ptr<func<void(i8)>>, i32) -> ()
+ llvm.return
+}
+
+// -----
+
+func.func @indirect_callee_return_mismatch(%callee : !llvm.ptr<func<i8()>>) {
+ // expected-error at +1 {{'llvm.call' op result type mismatch: 'i32' != 'i8'}}
+ "llvm.call"(%callee) : (!llvm.ptr<func<i8()>>) -> (i32)
+ llvm.return
+}
+
+// -----
+
+func.func @atomicrmw_mismatched_operands(%f32_ptr : !llvm.ptr<f32>, %i32 : i32) {
+ // expected-error at +1 {{expected LLVM IR element type for operand #0 to match type for operand #1}}
+ %0 = "llvm.atomicrmw"(%f32_ptr, %i32) {bin_op=11, ordering=1} : (!llvm.ptr<f32>, i32) -> i32
+ llvm.return
+}
+
+// -----
+
+func.func @cmpxchg_expected_ptr(%f32 : f32) {
+ // expected-error at +1 {{op operand #0 must be LLVM pointer to integer or LLVM pointer type}}
+ %0 = "llvm.cmpxchg"(%f32, %f32, %f32) {success_ordering=2,failure_ordering=2} : (f32, f32, f32) -> !llvm.struct<(f32, i1)>
+ llvm.return
+}
+
+// -----
+
+func.func @cmpxchg_mismatched_operands(%i64_ptr : !llvm.ptr<i64>, %i32 : i32) {
+ // expected-error at +1 {{expected LLVM IR element type for operand #0 to match type for all other operands}}
+ %0 = "llvm.cmpxchg"(%i64_ptr, %i32, %i32) {success_ordering=2,failure_ordering=2} : (!llvm.ptr<i64>, i32, i32) -> !llvm.struct<(i32, i1)>
+ llvm.return
+}
+
+// -----
+
+llvm.func @foo(i32) -> i32
+llvm.func @__gxx_personality_v0(...) -> i32
+
+llvm.func @bad_landingpad(%arg0: !llvm.ptr<ptr<i8>>) -> i32 attributes { personality = @__gxx_personality_v0} {
+ %0 = llvm.mlir.constant(3 : i32) : i32
+ %1 = llvm.mlir.constant(2 : i32) : i32
+ %2 = llvm.invoke @foo(%1) to ^bb1 unwind ^bb2 : (i32) -> i32
+^bb1: // pred: ^bb0
+ llvm.return %1 : i32
+^bb2: // pred: ^bb0
+ // expected-error at +1 {{clause #0 is not a known constant - null, addressof, bitcast}}
+ %3 = llvm.landingpad cleanup (catch %1 : i32) (catch %arg0 : !llvm.ptr<ptr<i8>>) : !llvm.struct<(ptr<i8>, i32)>
+ llvm.return %0 : i32
+}
+
+// -----
+
+llvm.func @foo(i32) -> i32
+llvm.func @__gxx_personality_v0(...) -> i32
+
+llvm.func @caller(%arg0: i32) -> i32 attributes { personality = @__gxx_personality_v0} {
+ %0 = llvm.mlir.constant(1 : i32) : i32
+ %1 = llvm.alloca %0 x !llvm.ptr<i8> : (i32) -> !llvm.ptr<ptr<i8>>
+ // expected-note at +1 {{global addresses expected as operand to bitcast used in clauses for landingpad}}
+ %2 = llvm.bitcast %1 : !llvm.ptr<ptr<i8>> to !llvm.ptr<i8>
+ %3 = llvm.invoke @foo(%0) to ^bb1 unwind ^bb2 : (i32) -> i32
+^bb1: // pred: ^bb0
+ llvm.return %0 : i32
+^bb2: // pred: ^bb0
+ // expected-error at +1 {{constant clauses expected}}
+ %5 = llvm.landingpad (catch %2 : !llvm.ptr<i8>) : !llvm.struct<(ptr<i8>, i32)>
+ llvm.return %0 : i32
+}
+
+// -----
+
+llvm.func @foo(i32) -> i32
+llvm.func @__gxx_personality_v0(...) -> i32
+
+llvm.func @caller(%arg0: i32) -> i32 attributes { personality = @__gxx_personality_v0} {
+ %0 = llvm.mlir.constant(1 : i32) : i32
+ %1 = llvm.invoke @foo(%0) to ^bb1 unwind ^bb2 : (i32) -> i32
+^bb1: // pred: ^bb0
+ llvm.return %0 : i32
+^bb2: // pred: ^bb0
+ // expected-error at +1 {{landingpad instruction expects at least one clause or cleanup attribute}}
+ %2 = llvm.landingpad : !llvm.struct<(ptr<i8>, i32)>
+ llvm.return %0 : i32
+}
+
+// -----
+
+llvm.func @foo(i32) -> i32
+llvm.func @__gxx_personality_v0(...) -> i32
+
+llvm.func @caller(%arg0: i32) -> i32 attributes { personality = @__gxx_personality_v0 } {
+ %0 = llvm.mlir.constant(1 : i32) : i32
+ %1 = llvm.invoke @foo(%0) to ^bb1 unwind ^bb2 : (i32) -> i32
+^bb1: // pred: ^bb0
+ llvm.return %0 : i32
+^bb2: // pred: ^bb0
+ %2 = llvm.landingpad cleanup : !llvm.struct<(ptr<i8>, i32)>
+ // expected-error at +1 {{'llvm.resume' op expects landingpad value as operand}}
+ llvm.resume %0 : i32
+}
+
+// -----
+
+llvm.func @foo(i32) -> i32
+
+llvm.func @caller(%arg0: i32) -> i32 {
+ %0 = llvm.mlir.constant(1 : i32) : i32
+ %1 = llvm.invoke @foo(%0) to ^bb1 unwind ^bb2 : (i32) -> i32
+^bb1: // pred: ^bb0
+ llvm.return %0 : i32
+^bb2: // pred: ^bb0
+ // expected-error at +1 {{llvm.landingpad needs to be in a function with a personality}}
+ %2 = llvm.landingpad cleanup : !llvm.struct<(ptr<i8>, i32)>
+ llvm.resume %2 : !llvm.struct<(ptr<i8>, i32)>
+}
+
+// -----
+
+llvm.func @wmmaLoadOp_invalid_mem_space(%arg0: !llvm.ptr<5>, %arg1: i32) {
+ // expected-error at +1 {{'nvvm.wmma.load' op expected source pointer in memory space 0, 1, 3}}
+ %0 = nvvm.wmma.load %arg0, %arg1
+ {eltype = #nvvm.mma_type<f16>, frag = #nvvm.mma_frag<a>, k = 16 : i32, layout = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}
+ : (!llvm.ptr<5>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
+ llvm.return
+}
+
+// -----
+
+llvm.func @wmmaLoadOp_invalid_AOp(%arg0: !llvm.ptr<3>, %arg1: i32) {
+ // expected-error at +1 {{'nvvm.wmma.load' op expected destination type is a structure of 8 elements of type 'vector<2xf16>'}}
+ %0 = nvvm.wmma.load %arg0, %arg1
+ {eltype = #nvvm.mma_type<f16>, frag = #nvvm.mma_frag<a>, k = 16 : i32, layout = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}
+ : (!llvm.ptr<3>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
+ llvm.return
+}
+
+// -----
+
+llvm.func @wmmaLoadOp_invalid_BOp(%arg0: !llvm.ptr<3>, %arg1: i32) {
+ // expected-error at +1 {{'nvvm.wmma.load' op expected destination type is a structure of 8 elements of type 'vector<2xf16>'}}
+ %0 = nvvm.wmma.load %arg0, %arg1
+ {eltype = #nvvm.mma_type<f16>, frag = #nvvm.mma_frag<b>, k = 16 : i32, layout = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}
+ : (!llvm.ptr<3>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
+
+ llvm.return
+}
+
+// -----
+
+llvm.func @wmmaLoadOp_invalid_COp(%arg0: !llvm.ptr<3>, %arg1: i32) {
+ // expected-error at +1 {{'nvvm.wmma.load' op expected destination type is a structure of 4 elements of type 'vector<2xf16>'}}
+ %0 = nvvm.wmma.load %arg0, %arg1
+ {eltype = #nvvm.mma_type<f16>, frag = #nvvm.mma_frag<c>, k = 16 : i32, layout = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}
+ : (!llvm.ptr<3>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>)>
+
+ llvm.return
+}
+
+// -----
+
+llvm.func @wmmaStoreOp_invalid_mem_space(%arg0: !llvm.ptr<5>, %arg1: i32,
+ %arg2: vector<2 x f16>, %arg3: vector<2 x f16>,
+ %arg4: vector<2 x f16>, %arg5: vector<2 xf16>) {
+ // expected-error at +1 {{'nvvm.wmma.store' op expected operands to be a source pointer in memory space 0, 1, 3}}
+ nvvm.wmma.store %arg0, %arg1, %arg2, %arg3, %arg4, %arg5
+ {eltype = #nvvm.mma_type<f16>, k = 16 : i32, layout = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}
+ : !llvm.ptr<5>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>
+ llvm.return
+}
+
+// -----
+
+llvm.func @wmmald_matrix(%arg0: !llvm.ptr<i32>) {
+ // expected-error at +1 {{'nvvm.ldmatrix' op expected source pointer in memory space 3}}
+ %l = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<i32>) -> i32
+ llvm.return
+}
+
+// -----
+
+llvm.func @wmmald_matrix(%arg0: !llvm.ptr<i32, 3>) {
+ // expected-error at +1 {{'nvvm.ldmatrix' op expected num attribute to be 1, 2 or 4}}
+ %l = nvvm.ldmatrix %arg0 {num = 3 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<i32, 3>) -> i32
+ llvm.return
+}
+
+// -----
+
+llvm.func @wmmald_matrix(%arg0: !llvm.ptr<i32, 3>) {
+ // expected-error at +1 {{'nvvm.ldmatrix' op expected destination type is i32}}
+ %l = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<i32, 3>) -> !llvm.struct<(i32)>
+ llvm.return
+}
+
+// -----
+
+llvm.func @wmmald_matrix(%arg0: !llvm.ptr<i32, 3>) {
+ // expected-error at +1 {{'nvvm.ldmatrix' op expected destination type is a structure of 4 elements of type i32}}
+ %l = nvvm.ldmatrix %arg0 {num = 4 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<i32, 3>) -> !llvm.struct<(i32, i32)>
+ llvm.return
+}
+
+// -----
+
+func.func @cp_async(%arg0: !llvm.ptr<i8, 3>, %arg1: !llvm.ptr<i8, 1>) {
+ // expected-error @below {{expected byte size to be either 4, 8 or 16.}}
+ nvvm.cp.async.shared.global %arg0, %arg1, 32 : !llvm.ptr<i8, 3>, !llvm.ptr<i8, 1>
+ return
+}
+
+// -----
+
+func.func @cp_async(%arg0: !llvm.ptr<i8, 3>, %arg1: !llvm.ptr<i8, 1>) {
+ // expected-error @below {{bypass l1 is only support for 16 bytes copy.}}
+ nvvm.cp.async.shared.global %arg0, %arg1, 8 {bypass_l1} : !llvm.ptr<i8, 3>, !llvm.ptr<i8, 1>
+ return
+}
+
+// -----
+
+func.func @gep_struct_variable(%arg0: !llvm.ptr<struct<(i32)>>, %arg1: i32, %arg2: i32) {
+ // expected-error @below {{op expected index 1 indexing a struct to be constant}}
+ llvm.getelementptr %arg0[%arg1, %arg1] : (!llvm.ptr<struct<(i32)>>, i32, i32) -> !llvm.ptr<i32>
+ return
+}
+
+// -----
+
+func.func @gep_out_of_bounds(%ptr: !llvm.ptr<struct<(i32, struct<(i32, f32)>)>>, %idx: i64) {
+ // expected-error @below {{index 2 indexing a struct is out of bounds}}
+ llvm.getelementptr %ptr[%idx, 1, 3] : (!llvm.ptr<struct<(i32, struct<(i32, f32)>)>>, i64) -> !llvm.ptr<i32>
+ return
+}
diff --git a/mlir/test/Dialect/LLVMIR/invalid.mlir b/mlir/test/Dialect/LLVMIR/invalid.mlir
index 3e019144a199..c3af84e55b88 100644
--- a/mlir/test/Dialect/LLVMIR/invalid.mlir
+++ b/mlir/test/Dialect/LLVMIR/invalid.mlir
@@ -64,7 +64,7 @@ func.func @alloca_missing_input_result_type(%size : i64) {
func.func @alloca_missing_input_type() {
// expected-error at +1 {{expected trailing function type with one argument and one result}}
- llvm.alloca %size x i32 : () -> (!llvm.ptr<i32>)
+ llvm.alloca %size x i32 : () -> (!llvm.ptr)
}
// -----
@@ -78,14 +78,14 @@ func.func @alloca_missing_result_type() {
func.func @alloca_non_function_type() {
// expected-error at +1 {{expected trailing function type with one argument and one result}}
- llvm.alloca %size x i32 : !llvm.ptr<i32>
+ llvm.alloca %size x i32 : !llvm.ptr
}
// -----
func.func @alloca_non_integer_alignment() {
// expected-error at +1 {{expected integer alignment}}
- llvm.alloca %size x i32 {alignment = 3.0} : !llvm.ptr<i32>
+ llvm.alloca %size x i32 {alignment = 3.0} : !llvm.ptr
}
// -----
@@ -97,44 +97,37 @@ func.func @alloca_opaque_ptr_no_type(%sz : i64) {
// -----
-func.func @alloca_ptr_type_attr_non_opaque_ptr(%sz : i64) {
- // expected-error at below {{unexpected 'elem_type' attribute when non-opaque pointer type is used}}
- "llvm.alloca"(%sz) { elem_type = i32 } : (i64) -> !llvm.ptr<i32>
-}
-
-// -----
-
-func.func @gep_missing_input_result_type(%pos : i64, %base : !llvm.ptr<f32>) {
+func.func @gep_missing_input_result_type(%pos : i64, %base : !llvm.ptr) {
// expected-error at +1 {{2 operands present, but expected 0}}
llvm.getelementptr %base[%pos] : () -> ()
}
// -----
-func.func @gep_missing_input_type(%pos : i64, %base : !llvm.ptr<f32>) {
+func.func @gep_missing_input_type(%pos : i64, %base : !llvm.ptr) {
// expected-error at +1 {{2 operands present, but expected 0}}
- llvm.getelementptr %base[%pos] : () -> (!llvm.ptr<f32>)
+ llvm.getelementptr %base[%pos] : () -> (!llvm.ptr)
}
// -----
-func.func @gep_missing_result_type(%pos : i64, %base : !llvm.ptr<f32>) {
+func.func @gep_missing_result_type(%pos : i64, %base : !llvm.ptr) {
// expected-error at +1 {{op requires one result}}
- llvm.getelementptr %base[%pos] : (!llvm.ptr<f32>, i64) -> ()
+ llvm.getelementptr %base[%pos] : (!llvm.ptr, i64) -> ()
}
// -----
-func.func @gep_non_function_type(%pos : i64, %base : !llvm.ptr<f32>) {
+func.func @gep_non_function_type(%pos : i64, %base : !llvm.ptr) {
// expected-error at +1 {{invalid kind of type specified}}
- llvm.getelementptr %base[%pos] : !llvm.ptr<f32>
+ llvm.getelementptr %base[%pos] : !llvm.ptr
}
// -----
-func.func @gep_too_few_dynamic(%base : !llvm.ptr<f32>) {
+func.func @gep_too_few_dynamic(%base : !llvm.ptr) {
// expected-error at +1 {{expected as many dynamic indices as specified in 'rawConstantIndices'}}
- %1 = "llvm.getelementptr"(%base) {rawConstantIndices = array<i32: -2147483648>} : (!llvm.ptr<f32>) -> !llvm.ptr<f32>
+ %1 = "llvm.getelementptr"(%base) {elem_type = f32, rawConstantIndices = array<i32: -2147483648>} : (!llvm.ptr) -> !llvm.ptr
}
// -----
@@ -302,14 +295,6 @@ func.func @call_unknown_symbol() {
// -----
-func.func @call_variadic(%callee : !llvm.ptr<func<i8 (i8, ...)>>, %arg : i8) {
- // expected-error at +1 {{indirect calls to variadic functions are not supported}}
- llvm.call %callee(%arg) : !llvm.ptr<func<i8 (i8, ...)>>, (i8) -> (i8)
- llvm.return
-}
-
-// -----
-
func.func private @standard_func_callee()
func.func @call_non_llvm() {
@@ -346,14 +331,6 @@ func.func @callee_arg_mismatch(%arg0 : i32) {
// -----
-func.func @indirect_callee_arg_mismatch(%arg0 : i32, %callee : !llvm.ptr<func<void(i8)>>) {
- // expected-error at +1 {{'llvm.call' op operand type mismatch for operand 0: 'i32' != 'i8'}}
- "llvm.call"(%callee, %arg0) : (!llvm.ptr<func<void(i8)>>, i32) -> ()
- llvm.return
-}
-
-// -----
-
llvm.func @callee_func() -> (i8)
func.func @callee_return_mismatch() {
@@ -364,14 +341,6 @@ func.func @callee_return_mismatch() {
// -----
-func.func @indirect_callee_return_mismatch(%callee : !llvm.ptr<func<i8()>>) {
- // expected-error at +1 {{'llvm.call' op result type mismatch: 'i32' != 'i8'}}
- "llvm.call"(%callee) : (!llvm.ptr<func<i8()>>) -> (i32)
- llvm.return
-}
-
-// -----
-
func.func @call_too_many_results(%callee : !llvm.ptr) {
// expected-error at +1 {{expected function with 0 or 1 result}}
llvm.call %callee() : !llvm.ptr, () -> (i32, i32)
@@ -406,14 +375,14 @@ llvm.func @func_result_mismatch(%arg0: f32) -> i32 {
func.func @constant_wrong_type() {
// expected-error at +1 {{only supports integer, float, string or elements attributes}}
- llvm.mlir.constant(@constant_wrong_type) : !llvm.ptr<func<void ()>>
+ llvm.mlir.constant(@constant_wrong_type) : !llvm.ptr
}
// -----
func.func @constant_wrong_type_string() {
// expected-error at below {{expected array type of 3 i8 elements for the string constant}}
- llvm.mlir.constant("foo") : !llvm.ptr<i8>
+ llvm.mlir.constant("foo") : !llvm.ptr
}
// -----
@@ -671,47 +640,39 @@ func.func @atomicrmw_expected_ptr(%f32 : f32) {
// -----
-func.func @atomicrmw_mismatched_operands(%f32_ptr : !llvm.ptr<f32>, %i32 : i32) {
- // expected-error at +1 {{expected LLVM IR element type for operand #0 to match type for operand #1}}
- %0 = "llvm.atomicrmw"(%f32_ptr, %i32) {bin_op=11, ordering=1} : (!llvm.ptr<f32>, i32) -> i32
- llvm.return
-}
-
-// -----
-
-func.func @atomicrmw_mismatched_operands(%f32_ptr : !llvm.ptr<f32>, %f32 : f32) {
+func.func @atomicrmw_mismatched_operands(%f32_ptr : !llvm.ptr, %f32 : f32) {
// expected-error at +1 {{op failed to verify that result #0 and operand #1 have the same type}}
- %0 = "llvm.atomicrmw"(%f32_ptr, %f32) {bin_op=11, ordering=1} : (!llvm.ptr<f32>, f32) -> i32
+ %0 = "llvm.atomicrmw"(%f32_ptr, %f32) {bin_op=11, ordering=1} : (!llvm.ptr, f32) -> i32
llvm.return
}
// -----
-func.func @atomicrmw_expected_float(%i32_ptr : !llvm.ptr<i32>, %i32 : i32) {
+func.func @atomicrmw_expected_float(%i32_ptr : !llvm.ptr, %i32 : i32) {
// expected-error at +1 {{expected LLVM IR floating point type}}
- %0 = llvm.atomicrmw fadd %i32_ptr, %i32 unordered : !llvm.ptr<i32>, i32
+ %0 = llvm.atomicrmw fadd %i32_ptr, %i32 unordered : !llvm.ptr, i32
llvm.return
}
// -----
-func.func @atomicrmw_unexpected_xchg_type(%i1_ptr : !llvm.ptr<i1>, %i1 : i1) {
+func.func @atomicrmw_unexpected_xchg_type(%i1_ptr : !llvm.ptr, %i1 : i1) {
// expected-error at +1 {{unexpected LLVM IR type for 'xchg' bin_op}}
- %0 = llvm.atomicrmw xchg %i1_ptr, %i1 unordered : !llvm.ptr<i1>, i1
+ %0 = llvm.atomicrmw xchg %i1_ptr, %i1 unordered : !llvm.ptr, i1
llvm.return
}
// -----
-func.func @atomicrmw_expected_int(%f32_ptr : !llvm.ptr<f32>, %f32 : f32) {
+func.func @atomicrmw_expected_int(%f32_ptr : !llvm.ptr, %f32 : f32) {
// expected-error at +1 {{expected LLVM IR integer type}}
- %0 = llvm.atomicrmw max %f32_ptr, %f32 unordered : !llvm.ptr<f32>, f32
+ %0 = llvm.atomicrmw max %f32_ptr, %f32 unordered : !llvm.ptr, f32
llvm.return
}
// -----
-func.func @cmpxchg_expected_ptr(%f32_ptr : !llvm.ptr<f32>, %f32 : f32) {
+func.func @cmpxchg_expected_ptr(%f32 : f32) {
// expected-error at +1 {{op operand #0 must be LLVM pointer to integer or LLVM pointer type}}
%0 = "llvm.cmpxchg"(%f32, %f32, %f32) {success_ordering=2,failure_ordering=2} : (f32, f32, f32) -> !llvm.struct<(f32, i1)>
llvm.return
@@ -719,14 +680,6 @@ func.func @cmpxchg_expected_ptr(%f32_ptr : !llvm.ptr<f32>, %f32 : f32) {
// -----
-func.func @cmpxchg_mismatched_operands(%i64_ptr : !llvm.ptr<i64>, %i32 : i32) {
- // expected-error at +1 {{expected LLVM IR element type for operand #0 to match type for all other operands}}
- %0 = "llvm.cmpxchg"(%i64_ptr, %i32, %i32) {success_ordering=2,failure_ordering=2} : (!llvm.ptr<i64>, i32, i32) -> !llvm.struct<(i32, i1)>
- llvm.return
-}
-
-// -----
-
func.func @cmpxchg_mismatched_value_operands(%ptr : !llvm.ptr, %i32 : i32, %i64 : i64) {
// expected-error at +1 {{op failed to verify that operand #1 and operand #2 have the same type}}
%0 = "llvm.cmpxchg"(%ptr, %i32, %i64) {success_ordering=2,failure_ordering=2} : (!llvm.ptr, i32, i64) -> !llvm.struct<(i32, i1)>
@@ -743,41 +696,41 @@ func.func @cmpxchg_mismatched_result(%ptr : !llvm.ptr, %i64 : i64) {
// -----
-func.func @cmpxchg_unexpected_type(%i1_ptr : !llvm.ptr<i1>, %i1 : i1) {
+func.func @cmpxchg_unexpected_type(%i1_ptr : !llvm.ptr, %i1 : i1) {
// expected-error at +1 {{unexpected LLVM IR type}}
- %0 = llvm.cmpxchg %i1_ptr, %i1, %i1 monotonic monotonic : !llvm.ptr<i1>, i1
+ %0 = llvm.cmpxchg %i1_ptr, %i1, %i1 monotonic monotonic : !llvm.ptr, i1
llvm.return
}
// -----
-func.func @cmpxchg_at_least_monotonic_success(%i32_ptr : !llvm.ptr<i32>, %i32 : i32) {
+func.func @cmpxchg_at_least_monotonic_success(%i32_ptr : !llvm.ptr, %i32 : i32) {
// expected-error at +1 {{ordering must be at least 'monotonic'}}
- %0 = llvm.cmpxchg %i32_ptr, %i32, %i32 unordered monotonic : !llvm.ptr<i32>, i32
+ %0 = llvm.cmpxchg %i32_ptr, %i32, %i32 unordered monotonic : !llvm.ptr, i32
llvm.return
}
// -----
-func.func @cmpxchg_at_least_monotonic_failure(%i32_ptr : !llvm.ptr<i32>, %i32 : i32) {
+func.func @cmpxchg_at_least_monotonic_failure(%i32_ptr : !llvm.ptr, %i32 : i32) {
// expected-error at +1 {{ordering must be at least 'monotonic'}}
- %0 = llvm.cmpxchg %i32_ptr, %i32, %i32 monotonic unordered : !llvm.ptr<i32>, i32
+ %0 = llvm.cmpxchg %i32_ptr, %i32, %i32 monotonic unordered : !llvm.ptr, i32
llvm.return
}
// -----
-func.func @cmpxchg_failure_release(%i32_ptr : !llvm.ptr<i32>, %i32 : i32) {
+func.func @cmpxchg_failure_release(%i32_ptr : !llvm.ptr, %i32 : i32) {
// expected-error at +1 {{failure ordering cannot be 'release' or 'acq_rel'}}
- %0 = llvm.cmpxchg %i32_ptr, %i32, %i32 acq_rel release : !llvm.ptr<i32>, i32
+ %0 = llvm.cmpxchg %i32_ptr, %i32, %i32 acq_rel release : !llvm.ptr, i32
llvm.return
}
// -----
-func.func @cmpxchg_failure_acq_rel(%i32_ptr : !llvm.ptr<i32>, %i32 : i32) {
+func.func @cmpxchg_failure_acq_rel(%i32_ptr : !llvm.ptr, %i32 : i32) {
// expected-error at +1 {{failure ordering cannot be 'release' or 'acq_rel'}}
- %0 = llvm.cmpxchg %i32_ptr, %i32, %i32 acq_rel acq_rel : !llvm.ptr<i32>, i32
+ %0 = llvm.cmpxchg %i32_ptr, %i32, %i32 acq_rel acq_rel : !llvm.ptr, i32
llvm.return
}
@@ -786,7 +739,7 @@ func.func @cmpxchg_failure_acq_rel(%i32_ptr : !llvm.ptr<i32>, %i32 : i32) {
llvm.func @foo(i32) -> i32
llvm.func @__gxx_personality_v0(...) -> i32
-llvm.func @bad_landingpad(%arg0: !llvm.ptr<ptr<i8>>) -> i32 attributes { personality = @__gxx_personality_v0} {
+llvm.func @bad_landingpad(%arg0: !llvm.ptr) -> i32 attributes { personality = @__gxx_personality_v0} {
%0 = llvm.mlir.constant(3 : i32) : i32
%1 = llvm.mlir.constant(2 : i32) : i32
%2 = llvm.invoke @foo(%1) to ^bb1 unwind ^bb2 : (i32) -> i32
@@ -794,7 +747,7 @@ llvm.func @bad_landingpad(%arg0: !llvm.ptr<ptr<i8>>) -> i32 attributes { persona
llvm.return %1 : i32
^bb2: // pred: ^bb0
// expected-error at +1 {{clause #0 is not a known constant - null, addressof, bitcast}}
- %3 = llvm.landingpad cleanup (catch %1 : i32) (catch %arg0 : !llvm.ptr<ptr<i8>>) : !llvm.struct<(ptr<i8>, i32)>
+ %3 = llvm.landingpad cleanup (catch %1 : i32) (catch %arg0 : !llvm.ptr) : !llvm.struct<(ptr, i32)>
llvm.return %0 : i32
}
@@ -805,15 +758,15 @@ llvm.func @__gxx_personality_v0(...) -> i32
llvm.func @caller(%arg0: i32) -> i32 attributes { personality = @__gxx_personality_v0} {
%0 = llvm.mlir.constant(1 : i32) : i32
- %1 = llvm.alloca %0 x !llvm.ptr<i8> : (i32) -> !llvm.ptr<ptr<i8>>
+ %1 = llvm.alloca %0 x !llvm.ptr : (i32) -> !llvm.ptr
// expected-note at +1 {{global addresses expected as operand to bitcast used in clauses for landingpad}}
- %2 = llvm.bitcast %1 : !llvm.ptr<ptr<i8>> to !llvm.ptr<i8>
+ %2 = llvm.bitcast %1 : !llvm.ptr to !llvm.ptr
%3 = llvm.invoke @foo(%0) to ^bb1 unwind ^bb2 : (i32) -> i32
^bb1: // pred: ^bb0
llvm.return %0 : i32
^bb2: // pred: ^bb0
// expected-error at +1 {{constant clauses expected}}
- %5 = llvm.landingpad (catch %2 : !llvm.ptr<i8>) : !llvm.struct<(ptr<i8>, i32)>
+ %5 = llvm.landingpad (catch %2 : !llvm.ptr) : !llvm.struct<(ptr, i32)>
llvm.return %0 : i32
}
@@ -829,7 +782,7 @@ llvm.func @caller(%arg0: i32) -> i32 attributes { personality = @__gxx_personali
llvm.return %0 : i32
^bb2: // pred: ^bb0
// expected-error at +1 {{landingpad instruction expects at least one clause or cleanup attribute}}
- %2 = llvm.landingpad : !llvm.struct<(ptr<i8>, i32)>
+ %2 = llvm.landingpad : !llvm.struct<(ptr, i32)>
llvm.return %0 : i32
}
@@ -844,7 +797,7 @@ llvm.func @caller(%arg0: i32) -> i32 attributes { personality = @__gxx_personali
^bb1: // pred: ^bb0
llvm.return %0 : i32
^bb2: // pred: ^bb0
- %2 = llvm.landingpad cleanup : !llvm.struct<(ptr<i8>, i32)>
+ %2 = llvm.landingpad cleanup : !llvm.struct<(ptr, i32)>
// expected-error at +1 {{'llvm.resume' op expects landingpad value as operand}}
llvm.resume %0 : i32
}
@@ -860,8 +813,8 @@ llvm.func @caller(%arg0: i32) -> i32 {
llvm.return %0 : i32
^bb2: // pred: ^bb0
// expected-error at +1 {{llvm.landingpad needs to be in a function with a personality}}
- %2 = llvm.landingpad cleanup : !llvm.struct<(ptr<i8>, i32)>
- llvm.resume %2 : !llvm.struct<(ptr<i8>, i32)>
+ %2 = llvm.landingpad cleanup : !llvm.struct<(ptr, i32)>
+ llvm.resume %2 : !llvm.struct<(ptr, i32)>
}
// -----
@@ -1056,55 +1009,55 @@ module {
// -----
-llvm.func @wmmaLoadOp_invalid_mem_space(%arg0: !llvm.ptr<i32, 5>, %arg1: i32) {
+llvm.func @wmmaLoadOp_invalid_mem_space(%arg0: !llvm.ptr<5>, %arg1: i32) {
// expected-error at +1 {{'nvvm.wmma.load' op expected source pointer in memory space 0, 1, 3}}
%0 = nvvm.wmma.load %arg0, %arg1
{eltype = #nvvm.mma_type<f16>, frag = #nvvm.mma_frag<a>, k = 16 : i32, layout = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}
- : (!llvm.ptr<i32, 5>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
+ : (!llvm.ptr<5>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
llvm.return
}
// -----
-llvm.func @wmmaLoadOp_invalid_AOp(%arg0: !llvm.ptr<i32, 3>, %arg1: i32) {
+llvm.func @wmmaLoadOp_invalid_AOp(%arg0: !llvm.ptr<3>, %arg1: i32) {
// expected-error at +1 {{'nvvm.wmma.load' op expected destination type is a structure of 8 elements of type 'vector<2xf16>'}}
%0 = nvvm.wmma.load %arg0, %arg1
{eltype = #nvvm.mma_type<f16>, frag = #nvvm.mma_frag<a>, k = 16 : i32, layout = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}
- : (!llvm.ptr<i32, 3>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
+ : (!llvm.ptr<3>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
llvm.return
}
// -----
-llvm.func @wmmaLoadOp_invalid_BOp(%arg0: !llvm.ptr<i32, 3>, %arg1: i32) {
+llvm.func @wmmaLoadOp_invalid_BOp(%arg0: !llvm.ptr<3>, %arg1: i32) {
// expected-error at +1 {{'nvvm.wmma.load' op expected destination type is a structure of 8 elements of type 'vector<2xf16>'}}
%0 = nvvm.wmma.load %arg0, %arg1
{eltype = #nvvm.mma_type<f16>, frag = #nvvm.mma_frag<b>, k = 16 : i32, layout = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}
- : (!llvm.ptr<i32, 3>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
+ : (!llvm.ptr<3>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>
llvm.return
}
// -----
-llvm.func @wmmaLoadOp_invalid_COp(%arg0: !llvm.ptr<i32, 3>, %arg1: i32) {
+llvm.func @wmmaLoadOp_invalid_COp(%arg0: !llvm.ptr<3>, %arg1: i32) {
// expected-error at +1 {{'nvvm.wmma.load' op expected destination type is a structure of 4 elements of type 'vector<2xf16>'}}
%0 = nvvm.wmma.load %arg0, %arg1
{eltype = #nvvm.mma_type<f16>, frag = #nvvm.mma_frag<c>, k = 16 : i32, layout = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}
- : (!llvm.ptr<i32, 3>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>)>
+ : (!llvm.ptr<3>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>)>
llvm.return
}
// -----
-llvm.func @wmmaStoreOp_invalid_mem_space(%arg0: !llvm.ptr<i32, 5>, %arg1: i32,
+llvm.func @wmmaStoreOp_invalid_mem_space(%arg0: !llvm.ptr<5>, %arg1: i32,
%arg2: vector<2 x f16>, %arg3: vector<2 x f16>,
%arg4: vector<2 x f16>, %arg5: vector<2 xf16>) {
// expected-error at +1 {{'nvvm.wmma.store' op expected operands to be a source pointer in memory space 0, 1, 3}}
nvvm.wmma.store %arg0, %arg1, %arg2, %arg3, %arg4, %arg5
{eltype = #nvvm.mma_type<f16>, k = 16 : i32, layout = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}
- : !llvm.ptr<i32, 5>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>
+ : !llvm.ptr<5>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>
llvm.return
}
@@ -1208,33 +1161,33 @@ llvm.func @gpu_wmma_mma_op_invalid_result(%arg0: vector<2 x f16>, %arg1: vector<
// -----
-llvm.func @wmmald_matrix(%arg0: !llvm.ptr<i32>) {
+llvm.func @wmmald_matrix(%arg0: !llvm.ptr) {
// expected-error at +1 {{'nvvm.ldmatrix' op expected source pointer in memory space 3}}
- %l = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<i32>) -> i32
+ %l = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr) -> i32
llvm.return
}
// -----
-llvm.func @wmmald_matrix(%arg0: !llvm.ptr<i32, 3>) {
+llvm.func @wmmald_matrix(%arg0: !llvm.ptr<3>) {
// expected-error at +1 {{'nvvm.ldmatrix' op expected num attribute to be 1, 2 or 4}}
- %l = nvvm.ldmatrix %arg0 {num = 3 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<i32, 3>) -> i32
+ %l = nvvm.ldmatrix %arg0 {num = 3 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<3>) -> i32
llvm.return
}
// -----
-llvm.func @wmmald_matrix(%arg0: !llvm.ptr<i32, 3>) {
+llvm.func @wmmald_matrix(%arg0: !llvm.ptr<3>) {
// expected-error at +1 {{'nvvm.ldmatrix' op expected destination type is i32}}
- %l = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<i32, 3>) -> !llvm.struct<(i32)>
+ %l = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<3>) -> !llvm.struct<(i32)>
llvm.return
}
// -----
-llvm.func @wmmald_matrix(%arg0: !llvm.ptr<i32, 3>) {
+llvm.func @wmmald_matrix(%arg0: !llvm.ptr<3>) {
// expected-error at +1 {{'nvvm.ldmatrix' op expected destination type is a structure of 4 elements of type i32}}
- %l = nvvm.ldmatrix %arg0 {num = 4 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<i32, 3>) -> !llvm.struct<(i32, i32)>
+ %l = nvvm.ldmatrix %arg0 {num = 4 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32)>
llvm.return
}
@@ -1278,33 +1231,33 @@ func.func @bitcast(%arg0: vector<2x3xf32>) {
// -----
-func.func @cp_async(%arg0: !llvm.ptr<i8, 3>, %arg1: !llvm.ptr<i8, 1>) {
+func.func @cp_async(%arg0: !llvm.ptr<3>, %arg1: !llvm.ptr<1>) {
// expected-error @below {{expected byte size to be either 4, 8 or 16.}}
- nvvm.cp.async.shared.global %arg0, %arg1, 32 : !llvm.ptr<i8, 3>, !llvm.ptr<i8, 1>
+ nvvm.cp.async.shared.global %arg0, %arg1, 32 : !llvm.ptr<3>, !llvm.ptr<1>
return
}
// -----
-func.func @cp_async(%arg0: !llvm.ptr<i8, 3>, %arg1: !llvm.ptr<i8, 1>) {
+func.func @cp_async(%arg0: !llvm.ptr<3>, %arg1: !llvm.ptr<1>) {
// expected-error @below {{bypass l1 is only support for 16 bytes copy.}}
- nvvm.cp.async.shared.global %arg0, %arg1, 8 {bypass_l1} : !llvm.ptr<i8, 3>, !llvm.ptr<i8, 1>
+ nvvm.cp.async.shared.global %arg0, %arg1, 8 {bypass_l1} : !llvm.ptr<3>, !llvm.ptr<1>
return
}
// -----
-func.func @gep_struct_variable(%arg0: !llvm.ptr<struct<(i32)>>, %arg1: i32, %arg2: i32) {
+func.func @gep_struct_variable(%arg0: !llvm.ptr, %arg1: i32, %arg2: i32) {
// expected-error @below {{op expected index 1 indexing a struct to be constant}}
- llvm.getelementptr %arg0[%arg1, %arg1] : (!llvm.ptr<struct<(i32)>>, i32, i32) -> !llvm.ptr<i32>
+ llvm.getelementptr %arg0[%arg1, %arg1] : (!llvm.ptr, i32, i32) -> !llvm.ptr, !llvm.struct<(i32)>
return
}
// -----
-func.func @gep_out_of_bounds(%ptr: !llvm.ptr<struct<(i32, struct<(i32, f32)>)>>, %idx: i64) {
+func.func @gep_out_of_bounds(%ptr: !llvm.ptr, %idx: i64) {
// expected-error @below {{index 2 indexing a struct is out of bounds}}
- llvm.getelementptr %ptr[%idx, 1, 3] : (!llvm.ptr<struct<(i32, struct<(i32, f32)>)>>, i64) -> !llvm.ptr<i32>
+ llvm.getelementptr %ptr[%idx, 1, 3] : (!llvm.ptr, i64) -> !llvm.ptr, !llvm.struct<(i32, struct<(i32, f32)>)>
return
}
@@ -1321,8 +1274,8 @@ func.func @non_splat_shuffle_on_scalable_vector(%arg0: vector<[4]xf32>) {
llvm.mlir.global internal @side_effecting_global() : !llvm.struct<(i8)> {
%0 = llvm.mlir.constant(1 : i64) : i64
// expected-error at below {{ops with side effects not allowed in global initializers}}
- %1 = llvm.alloca %0 x !llvm.struct<(i8)> : (i64) -> !llvm.ptr<struct<(i8)>>
- %2 = llvm.load %1 : !llvm.ptr<struct<(i8)>>
+ %1 = llvm.alloca %0 x !llvm.struct<(i8)> : (i64) -> !llvm.ptr
+ %2 = llvm.load %1 : !llvm.ptr -> !llvm.struct<(i8)>
llvm.return %2 : !llvm.struct<(i8)>
}
diff --git a/mlir/test/Dialect/LLVMIR/layout-typed-pointers.mlir b/mlir/test/Dialect/LLVMIR/layout-typed-pointers.mlir
new file mode 100644
index 000000000000..5cf1ed03e64c
--- /dev/null
+++ b/mlir/test/Dialect/LLVMIR/layout-typed-pointers.mlir
@@ -0,0 +1,145 @@
+// RUN: mlir-opt --test-data-layout-query --split-input-file --verify-diagnostics %s | FileCheck %s
+
+module {
+ // CHECK: @no_spec
+ func.func @no_spec() {
+ // CHECK: alignment = 8
+ // CHECK: alloca_memory_space = 0
+ // CHECK: bitsize = 64
+ // CHECK: preferred = 8
+ // CHECK: size = 8
+ "test.data_layout_query"() : () -> !llvm.ptr<i8>
+ // CHECK: alignment = 8
+ // CHECK: alloca_memory_space = 0
+ // CHECK: bitsize = 64
+ // CHECK: preferred = 8
+ // CHECK: size = 8
+ "test.data_layout_query"() : () -> !llvm.ptr<i32>
+ // CHECK: alignment = 8
+ // CHECK: alloca_memory_space = 0
+ // CHECK: bitsize = 64
+ // CHECK: preferred = 8
+ // CHECK: size = 8
+ "test.data_layout_query"() : () -> !llvm.ptr<bf16>
+ // CHECK: alignment = 8
+ // CHECK: alloca_memory_space = 0
+ // CHECK: bitsize = 64
+ // CHECK: preferred = 8
+ // CHECK: size = 8
+ "test.data_layout_query"() : () -> !llvm.ptr<!llvm.ptr<i8>>
+ // CHECK: alignment = 8
+ // CHECK: alloca_memory_space = 0
+ // CHECK: bitsize = 64
+ // CHECK: preferred = 8
+ // CHECK: size = 8
+ "test.data_layout_query"() : () -> !llvm.ptr<i8, 3>
+ // CHECK: alignment = 8
+ // CHECK: alloca_memory_space = 0
+ // CHECK: bitsize = 64
+ // CHECK: preferred = 8
+ // CHECK: size = 8
+ "test.data_layout_query"() : () -> !llvm.ptr<i8, 5>
+ // CHECK: alignment = 8
+ // CHECK: alloca_memory_space = 0
+ // CHECK: bitsize = 64
+ // CHECK: preferred = 8
+ // CHECK: size = 8
+ "test.data_layout_query"() : () -> !llvm.ptr<5>
+ return
+ }
+}
+
+// -----
+
+module attributes { dlti.dl_spec = #dlti.dl_spec<
+ #dlti.dl_entry<!llvm.ptr<i8>, dense<[32, 32, 64]> : vector<3xi32>>,
+ #dlti.dl_entry<!llvm.ptr<i8, 5>, dense<[64, 64, 64]> : vector<3xi32>>,
+ #dlti.dl_entry<!llvm.ptr<4>, dense<[32, 64, 64]> : vector<3xi32>>,
+ #dlti.dl_entry<"dlti.alloca_memory_space", 5 : ui32>
+>} {
+ // CHECK: @spec
+ func.func @spec() {
+ // CHECK: alignment = 4
+ // CHECK: alloca_memory_space = 5
+ // CHECK: bitsize = 32
+ // CHECK: preferred = 8
+ // CHECK: size = 4
+ "test.data_layout_query"() : () -> !llvm.ptr<i8>
+ // CHECK: alignment = 4
+ // CHECK: alloca_memory_space = 5
+ // CHECK: bitsize = 32
+ // CHECK: preferred = 8
+ // CHECK: size = 4
+ "test.data_layout_query"() : () -> !llvm.ptr<i32>
+ // CHECK: alignment = 4
+ // CHECK: alloca_memory_space = 5
+ // CHECK: bitsize = 32
+ // CHECK: preferred = 8
+ // CHECK: size = 4
+ "test.data_layout_query"() : () -> !llvm.ptr<bf16>
+ // CHECK: alignment = 4
+ // CHECK: alloca_memory_space = 5
+ // CHECK: bitsize = 32
+ // CHECK: preferred = 8
+ // CHECK: size = 4
+ "test.data_layout_query"() : () -> !llvm.ptr<!llvm.ptr<i8>>
+ // CHECK: alignment = 4
+ // CHECK: alloca_memory_space = 5
+ // CHECK: bitsize = 32
+ // CHECK: preferred = 8
+ // CHECK: size = 4
+ "test.data_layout_query"() : () -> !llvm.ptr<i8, 3>
+ // CHECK: alignment = 8
+ // CHECK: alloca_memory_space = 5
+ // CHECK: bitsize = 64
+ // CHECK: preferred = 8
+ // CHECK: size = 8
+ "test.data_layout_query"() : () -> !llvm.ptr<i8, 5>
+ // CHECK: alignment = 4
+ // CHECK: alloca_memory_space = 5
+ // CHECK: bitsize = 32
+ // CHECK: preferred = 8
+ // CHECK: size = 4
+ "test.data_layout_query"() : () -> !llvm.ptr<3>
+ // CHECK: alignment = 8
+ // CHECK: alloca_memory_space = 5
+ // CHECK: bitsize = 32
+ // CHECK: preferred = 8
+ // CHECK: size = 4
+ "test.data_layout_query"() : () -> !llvm.ptr<4>
+ return
+ }
+}
+
+// -----
+
+// expected-error at below {{unexpected layout attribute for pointer to 'i32'}}
+module attributes { dlti.dl_spec = #dlti.dl_spec<
+ #dlti.dl_entry<!llvm.ptr<i32>, dense<[64, 64, 64]> : vector<3xi32>>
+>} {
+ func.func @pointer() {
+ return
+ }
+}
+
+// -----
+
+// expected-error at below {{expected layout attribute for '!llvm.ptr<i8>' to be a dense integer elements attribute with 3 or 4 elements}}
+module attributes { dlti.dl_spec = #dlti.dl_spec<
+ #dlti.dl_entry<!llvm.ptr<i8>, dense<[64.0, 64.0, 64.0]> : vector<3xf32>>
+>} {
+ func.func @pointer() {
+ return
+ }
+}
+
+// -----
+
+// expected-error at below {{preferred alignment is expected to be at least as large as ABI alignment}}
+module attributes { dlti.dl_spec = #dlti.dl_spec<
+ #dlti.dl_entry<!llvm.ptr<i8>, dense<[64, 64, 32]> : vector<3xi32>>
+>} {
+ func.func @pointer() {
+ return
+ }
+}
diff --git a/mlir/test/Dialect/LLVMIR/layout.mlir b/mlir/test/Dialect/LLVMIR/layout.mlir
index e5c8c0bd86db..d6e2013cc86c 100644
--- a/mlir/test/Dialect/LLVMIR/layout.mlir
+++ b/mlir/test/Dialect/LLVMIR/layout.mlir
@@ -3,42 +3,13 @@
module {
// CHECK: @no_spec
func.func @no_spec() {
+ "test.data_layout_query"() : () -> !llvm.ptr
// CHECK: alignment = 8
// CHECK: alloca_memory_space = 0
// CHECK: bitsize = 64
// CHECK: preferred = 8
// CHECK: size = 8
- "test.data_layout_query"() : () -> !llvm.ptr<i8>
- // CHECK: alignment = 8
- // CHECK: alloca_memory_space = 0
- // CHECK: bitsize = 64
- // CHECK: preferred = 8
- // CHECK: size = 8
- "test.data_layout_query"() : () -> !llvm.ptr<i32>
- // CHECK: alignment = 8
- // CHECK: alloca_memory_space = 0
- // CHECK: bitsize = 64
- // CHECK: preferred = 8
- // CHECK: size = 8
- "test.data_layout_query"() : () -> !llvm.ptr<bf16>
- // CHECK: alignment = 8
- // CHECK: alloca_memory_space = 0
- // CHECK: bitsize = 64
- // CHECK: preferred = 8
- // CHECK: size = 8
- "test.data_layout_query"() : () -> !llvm.ptr<!llvm.ptr<i8>>
- // CHECK: alignment = 8
- // CHECK: alloca_memory_space = 0
- // CHECK: bitsize = 64
- // CHECK: preferred = 8
- // CHECK: size = 8
- "test.data_layout_query"() : () -> !llvm.ptr<i8, 3>
- // CHECK: alignment = 8
- // CHECK: alloca_memory_space = 0
- // CHECK: bitsize = 64
- // CHECK: preferred = 8
- // CHECK: size = 8
- "test.data_layout_query"() : () -> !llvm.ptr<i8, 5>
+ "test.data_layout_query"() : () -> !llvm.ptr<3>
// CHECK: alignment = 8
// CHECK: alloca_memory_space = 0
// CHECK: bitsize = 64
@@ -52,8 +23,8 @@ module {
// -----
module attributes { dlti.dl_spec = #dlti.dl_spec<
- #dlti.dl_entry<!llvm.ptr<i8>, dense<[32, 32, 64]> : vector<3xi32>>,
- #dlti.dl_entry<!llvm.ptr<i8, 5>, dense<[64, 64, 64]> : vector<3xi32>>,
+ #dlti.dl_entry<!llvm.ptr, dense<[32, 32, 64]> : vector<3xi32>>,
+ #dlti.dl_entry<!llvm.ptr<5>, dense<[64, 64, 64]> : vector<3xi32>>,
#dlti.dl_entry<!llvm.ptr<4>, dense<[32, 64, 64]> : vector<3xi32>>,
#dlti.dl_entry<"dlti.alloca_memory_space", 5 : ui32>
>} {
@@ -64,37 +35,19 @@ module attributes { dlti.dl_spec = #dlti.dl_spec<
// CHECK: bitsize = 32
// CHECK: preferred = 8
// CHECK: size = 4
- "test.data_layout_query"() : () -> !llvm.ptr<i8>
+ "test.data_layout_query"() : () -> !llvm.ptr
// CHECK: alignment = 4
// CHECK: alloca_memory_space = 5
// CHECK: bitsize = 32
// CHECK: preferred = 8
// CHECK: size = 4
- "test.data_layout_query"() : () -> !llvm.ptr<i32>
- // CHECK: alignment = 4
- // CHECK: alloca_memory_space = 5
- // CHECK: bitsize = 32
- // CHECK: preferred = 8
- // CHECK: size = 4
- "test.data_layout_query"() : () -> !llvm.ptr<bf16>
- // CHECK: alignment = 4
- // CHECK: alloca_memory_space = 5
- // CHECK: bitsize = 32
- // CHECK: preferred = 8
- // CHECK: size = 4
- "test.data_layout_query"() : () -> !llvm.ptr<!llvm.ptr<i8>>
- // CHECK: alignment = 4
- // CHECK: alloca_memory_space = 5
- // CHECK: bitsize = 32
- // CHECK: preferred = 8
- // CHECK: size = 4
- "test.data_layout_query"() : () -> !llvm.ptr<i8, 3>
+ "test.data_layout_query"() : () -> !llvm.ptr<3>
// CHECK: alignment = 8
// CHECK: alloca_memory_space = 5
// CHECK: bitsize = 64
// CHECK: preferred = 8
// CHECK: size = 8
- "test.data_layout_query"() : () -> !llvm.ptr<i8, 5>
+ "test.data_layout_query"() : () -> !llvm.ptr<5>
// CHECK: alignment = 4
// CHECK: alloca_memory_space = 5
// CHECK: bitsize = 32
@@ -113,20 +66,9 @@ module attributes { dlti.dl_spec = #dlti.dl_spec<
// -----
-// expected-error at below {{unexpected layout attribute for pointer to 'i32'}}
-module attributes { dlti.dl_spec = #dlti.dl_spec<
- #dlti.dl_entry<!llvm.ptr<i32>, dense<[64, 64, 64]> : vector<3xi32>>
->} {
- func.func @pointer() {
- return
- }
-}
-
-// -----
-
-// expected-error at below {{expected layout attribute for '!llvm.ptr<i8>' to be a dense integer elements attribute with 3 or 4 elements}}
+// expected-error at below {{expected layout attribute for '!llvm.ptr' to be a dense integer elements attribute with 3 or 4 elements}}
module attributes { dlti.dl_spec = #dlti.dl_spec<
- #dlti.dl_entry<!llvm.ptr<i8>, dense<[64.0, 64.0, 64.0]> : vector<3xf32>>
+ #dlti.dl_entry<!llvm.ptr, dense<[64.0, 64.0, 64.0]> : vector<3xf32>>
>} {
func.func @pointer() {
return
@@ -137,7 +79,7 @@ module attributes { dlti.dl_spec = #dlti.dl_spec<
// expected-error at below {{preferred alignment is expected to be at least as large as ABI alignment}}
module attributes { dlti.dl_spec = #dlti.dl_spec<
- #dlti.dl_entry<!llvm.ptr<i8>, dense<[64, 64, 32]> : vector<3xi32>>
+ #dlti.dl_entry<!llvm.ptr, dense<[64, 64, 32]> : vector<3xi32>>
>} {
func.func @pointer() {
return
diff --git a/mlir/test/Dialect/LLVMIR/nvvm-typed-pointers.mlir b/mlir/test/Dialect/LLVMIR/nvvm-typed-pointers.mlir
new file mode 100644
index 000000000000..5fbadd1dc414
--- /dev/null
+++ b/mlir/test/Dialect/LLVMIR/nvvm-typed-pointers.mlir
@@ -0,0 +1,55 @@
+// RUN: mlir-opt %s -split-input-file -verify-diagnostics | FileCheck %s
+
+// CHECK-LABEL: @nvvm_wmma_load_tf32
+func.func @nvvm_wmma_load_tf32(%arg0: !llvm.ptr<i32>, %arg1 : i32) -> !llvm.struct<(i32, i32, i32, i32)> {
+ // CHECK: nvvm.wmma.load {{.*}} {eltype = #nvvm.mma_type<tf32>, frag = #nvvm.mma_frag<a>, k = 8 : i32, layout = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}
+ %0 = nvvm.wmma.load %arg0, %arg1
+ {eltype = #nvvm.mma_type<tf32>, frag = #nvvm.mma_frag<a>, k = 8 : i32, layout = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}
+ : (!llvm.ptr<i32>) -> !llvm.struct<(i32, i32, i32, i32)>
+ llvm.return %0 : !llvm.struct<(i32, i32, i32, i32)>
+}
+
+// CHECK-LABEL: @cp_async
+llvm.func @cp_async(%arg0: !llvm.ptr<i8, 3>, %arg1: !llvm.ptr<i8, 1>) {
+// CHECK: nvvm.cp.async.shared.global %{{.*}}, %{{.*}}, 16
+ nvvm.cp.async.shared.global %arg0, %arg1, 16 : !llvm.ptr<i8, 3>, !llvm.ptr<i8, 1>
+// CHECK: nvvm.cp.async.shared.global %{{.*}}, %{{.*}}, 16 {bypass_l1}
+ nvvm.cp.async.shared.global %arg0, %arg1, 16 {bypass_l1} : !llvm.ptr<i8, 3>, !llvm.ptr<i8, 1>
+// CHECK: nvvm.cp.async.commit.group
+ nvvm.cp.async.commit.group
+// CHECK: nvvm.cp.async.wait.group 0
+ nvvm.cp.async.wait.group 0
+ llvm.return
+}
+
+// CHECK-LABEL: llvm.func @ld_matrix
+llvm.func @ld_matrix(%arg0: !llvm.ptr<i32, 3>) {
+ // CHECK: nvvm.ldmatrix %{{.*}} {layout = #nvvm.mma_layout<row>, num = 1 : i32} : (!llvm.ptr<i32, 3>) -> i32
+ %l1 = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<i32, 3>) -> i32
+ // CHECK: nvvm.ldmatrix %{{.*}} {layout = #nvvm.mma_layout<row>, num = 2 : i32} : (!llvm.ptr<i32, 3>) -> !llvm.struct<(i32, i32)>
+ %l2 = nvvm.ldmatrix %arg0 {num = 2 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<i32, 3>) -> !llvm.struct<(i32, i32)>
+ // CHECK: nvvm.ldmatrix %{{.*}} {layout = #nvvm.mma_layout<row>, num = 4 : i32} : (!llvm.ptr<i32, 3>) -> !llvm.struct<(i32, i32, i32, i32)>
+ %l4 = nvvm.ldmatrix %arg0 {num = 4 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<i32, 3>) -> !llvm.struct<(i32, i32, i32, i32)>
+ llvm.return
+}
+
+// CHECK-LABEL: llvm.func @redux_sync
+llvm.func @redux_sync(%value : i32, %offset : i32) -> i32 {
+ // CHECK: nvvm.redux.sync add %{{.*}}
+ %r1 = nvvm.redux.sync add %value, %offset : i32 -> i32
+ // CHECK: nvvm.redux.sync max %{{.*}}
+ %r2 = nvvm.redux.sync max %value, %offset : i32 -> i32
+ // CHECK: nvvm.redux.sync min %{{.*}}
+ %r3 = nvvm.redux.sync min %value, %offset : i32 -> i32
+ // CHECK: nvvm.redux.sync umax %{{.*}}
+ %r5 = nvvm.redux.sync umax %value, %offset : i32 -> i32
+ // CHECK: nvvm.redux.sync umin %{{.*}}
+ %r6 = nvvm.redux.sync umin %value, %offset : i32 -> i32
+ // CHECK: nvvm.redux.sync and %{{.*}}
+ %r7 = nvvm.redux.sync and %value, %offset : i32 -> i32
+ // CHECK: nvvm.redux.sync or %{{.*}}
+ %r8 = nvvm.redux.sync or %value, %offset : i32 -> i32
+ // CHECK: nvvm.redux.sync xor %{{.*}}
+ %r9 = nvvm.redux.sync xor %value, %offset : i32 -> i32
+ llvm.return %r1 : i32
+}
diff --git a/mlir/test/Dialect/LLVMIR/nvvm.mlir b/mlir/test/Dialect/LLVMIR/nvvm.mlir
index 6596b8503d7a..c7c83d29638c 100644
--- a/mlir/test/Dialect/LLVMIR/nvvm.mlir
+++ b/mlir/test/Dialect/LLVMIR/nvvm.mlir
@@ -266,11 +266,11 @@ func.func @nvvm_mma_m16n8k32_s4_s4(%a0 : i32, %a1 : i32,
}
// CHECK-LABEL: @nvvm_wmma_load_tf32
-func.func @nvvm_wmma_load_tf32(%arg0: !llvm.ptr<i32>, %arg1 : i32) -> !llvm.struct<(i32, i32, i32, i32)> {
+func.func @nvvm_wmma_load_tf32(%arg0: !llvm.ptr, %arg1 : i32) -> !llvm.struct<(i32, i32, i32, i32)> {
// CHECK: nvvm.wmma.load {{.*}} {eltype = #nvvm.mma_type<tf32>, frag = #nvvm.mma_frag<a>, k = 8 : i32, layout = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}
%0 = nvvm.wmma.load %arg0, %arg1
{eltype = #nvvm.mma_type<tf32>, frag = #nvvm.mma_frag<a>, k = 8 : i32, layout = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}
- : (!llvm.ptr<i32>) -> !llvm.struct<(i32, i32, i32, i32)>
+ : (!llvm.ptr) -> !llvm.struct<(i32, i32, i32, i32)>
llvm.return %0 : !llvm.struct<(i32, i32, i32, i32)>
}
@@ -288,11 +288,11 @@ func.func @nvvm_wmma_mma(%0 : i32, %1 : i32, %2 : i32, %3 : i32, %4 : i32, %5 :
}
// CHECK-LABEL: @cp_async
-llvm.func @cp_async(%arg0: !llvm.ptr<i8, 3>, %arg1: !llvm.ptr<i8, 1>) {
+llvm.func @cp_async(%arg0: !llvm.ptr<3>, %arg1: !llvm.ptr<1>) {
// CHECK: nvvm.cp.async.shared.global %{{.*}}, %{{.*}}, 16
- nvvm.cp.async.shared.global %arg0, %arg1, 16 : !llvm.ptr<i8, 3>, !llvm.ptr<i8, 1>
+ nvvm.cp.async.shared.global %arg0, %arg1, 16 : !llvm.ptr<3>, !llvm.ptr<1>
// CHECK: nvvm.cp.async.shared.global %{{.*}}, %{{.*}}, 16 {bypass_l1}
- nvvm.cp.async.shared.global %arg0, %arg1, 16 {bypass_l1} : !llvm.ptr<i8, 3>, !llvm.ptr<i8, 1>
+ nvvm.cp.async.shared.global %arg0, %arg1, 16 {bypass_l1} : !llvm.ptr<3>, !llvm.ptr<1>
// CHECK: nvvm.cp.async.commit.group
nvvm.cp.async.commit.group
// CHECK: nvvm.cp.async.wait.group 0
@@ -301,18 +301,18 @@ llvm.func @cp_async(%arg0: !llvm.ptr<i8, 3>, %arg1: !llvm.ptr<i8, 1>) {
}
// CHECK-LABEL: llvm.func @ld_matrix
-llvm.func @ld_matrix(%arg0: !llvm.ptr<i32, 3>) {
- // CHECK: nvvm.ldmatrix %{{.*}} {layout = #nvvm.mma_layout<row>, num = 1 : i32} : (!llvm.ptr<i32, 3>) -> i32
- %l1 = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<i32, 3>) -> i32
- // CHECK: nvvm.ldmatrix %{{.*}} {layout = #nvvm.mma_layout<row>, num = 2 : i32} : (!llvm.ptr<i32, 3>) -> !llvm.struct<(i32, i32)>
- %l2 = nvvm.ldmatrix %arg0 {num = 2 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<i32, 3>) -> !llvm.struct<(i32, i32)>
- // CHECK: nvvm.ldmatrix %{{.*}} {layout = #nvvm.mma_layout<row>, num = 4 : i32} : (!llvm.ptr<i32, 3>) -> !llvm.struct<(i32, i32, i32, i32)>
- %l4 = nvvm.ldmatrix %arg0 {num = 4 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<i32, 3>) -> !llvm.struct<(i32, i32, i32, i32)>
+llvm.func @ld_matrix(%arg0: !llvm.ptr<3>) {
+ // CHECK: nvvm.ldmatrix %{{.*}} {layout = #nvvm.mma_layout<row>, num = 1 : i32} : (!llvm.ptr<3>) -> i32
+ %l1 = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<3>) -> i32
+ // CHECK: nvvm.ldmatrix %{{.*}} {layout = #nvvm.mma_layout<row>, num = 2 : i32} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32)>
+ %l2 = nvvm.ldmatrix %arg0 {num = 2 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32)>
+ // CHECK: nvvm.ldmatrix %{{.*}} {layout = #nvvm.mma_layout<row>, num = 4 : i32} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32, i32, i32)>
+ %l4 = nvvm.ldmatrix %arg0 {num = 4 : i32, layout = #nvvm.mma_layout<row>} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32, i32, i32)>
llvm.return
}
// CHECK-LABEL: llvm.func @redux_sync
-llvm.func @redux_sync(%value : i32, %offset : i32) -> i32 {
+llvm.func @redux_sync(%value : i32, %offset : i32) -> i32 {
// CHECK: nvvm.redux.sync add %{{.*}}
%r1 = nvvm.redux.sync add %value, %offset : i32 -> i32
// CHECK: nvvm.redux.sync max %{{.*}}
@@ -324,9 +324,9 @@ llvm.func @redux_sync(%value : i32, %offset : i32) -> i32 {
// CHECK: nvvm.redux.sync umin %{{.*}}
%r6 = nvvm.redux.sync umin %value, %offset : i32 -> i32
// CHECK: nvvm.redux.sync and %{{.*}}
- %r7 = nvvm.redux.sync and %value, %offset : i32 -> i32
+ %r7 = nvvm.redux.sync and %value, %offset : i32 -> i32
// CHECK: nvvm.redux.sync or %{{.*}}
- %r8 = nvvm.redux.sync or %value, %offset : i32 -> i32
+ %r8 = nvvm.redux.sync or %value, %offset : i32 -> i32
// CHECK: nvvm.redux.sync xor %{{.*}}
%r9 = nvvm.redux.sync xor %value, %offset : i32 -> i32
llvm.return %r1 : i32
diff --git a/mlir/test/Dialect/LLVMIR/parameter-attrs-invalid-typed-pointers.mlir b/mlir/test/Dialect/LLVMIR/parameter-attrs-invalid-typed-pointers.mlir
new file mode 100644
index 000000000000..65411ff41e28
--- /dev/null
+++ b/mlir/test/Dialect/LLVMIR/parameter-attrs-invalid-typed-pointers.mlir
@@ -0,0 +1,6 @@
+// RUN: mlir-opt %s -split-input-file -verify-diagnostics
+
+// Argument attributes
+
+// expected-error at below {{"llvm.sret" attribute attached to LLVM pointer argument of
diff erent type}}
+llvm.func @invalid_sret_attr_type(%0 : !llvm.ptr<f32> {llvm.sret = !llvm.struct<(i32)>})
diff --git a/mlir/test/Dialect/LLVMIR/parameter-attrs-invalid.mlir b/mlir/test/Dialect/LLVMIR/parameter-attrs-invalid.mlir
index 72bf45052ef1..d7ee6097b360 100644
--- a/mlir/test/Dialect/LLVMIR/parameter-attrs-invalid.mlir
+++ b/mlir/test/Dialect/LLVMIR/parameter-attrs-invalid.mlir
@@ -47,11 +47,6 @@ llvm.func @invalid_sret_arg_type(%0 : i32 {llvm.sret = !llvm.struct<(i32)>})
// -----
-// expected-error at below {{"llvm.sret" attribute attached to LLVM pointer argument of
diff erent type}}
-llvm.func @invalid_sret_attr_type(%0 : !llvm.ptr<f32> {llvm.sret = !llvm.struct<(i32)>})
-
-// -----
-
// expected-error at below {{"llvm.byval" attribute attached to non-pointer LLVM type}}
llvm.func @invalid_byval_arg_type(%0 : i32 {llvm.byval = !llvm.struct<(i32)>})
diff --git a/mlir/test/Dialect/LLVMIR/types-invalid-typed-pointers.mlir b/mlir/test/Dialect/LLVMIR/types-invalid-typed-pointers.mlir
new file mode 100644
index 000000000000..475fadede8fb
--- /dev/null
+++ b/mlir/test/Dialect/LLVMIR/types-invalid-typed-pointers.mlir
@@ -0,0 +1,42 @@
+// RUN: mlir-opt --allow-unregistered-dialect -split-input-file -verify-diagnostics %s
+
+func.func @void_pointer() {
+ // expected-error @+1 {{invalid pointer element type}}
+ "some.op"() : () -> !llvm.ptr<void>
+}
+
+// -----
+
+func.func @repeated_struct_name() {
+ "some.op"() : () -> !llvm.struct<"a", (ptr<struct<"a">>)>
+ // expected-error @+1 {{identified type already used with a
diff erent body}}
+ "some.op"() : () -> !llvm.struct<"a", (i32)>
+}
+
+// -----
+
+func.func @dynamic_vector() {
+ // expected-error @+1 {{expected '? x <integer> x <type>' or '<integer> x <type>'}}
+ "some.op"() : () -> !llvm.vec<? x ptr<f32>>
+}
+
+// -----
+
+func.func @dynamic_scalable_vector() {
+ // expected-error @+1 {{expected '? x <integer> x <type>' or '<integer> x <type>'}}
+ "some.op"() : () -> !llvm.vec<?x? x ptr<f32>>
+}
+
+// -----
+
+func.func @unscalable_vector() {
+ // expected-error @+1 {{expected '? x <integer> x <type>' or '<integer> x <type>'}}
+ "some.op"() : () -> !llvm.vec<4x4 x ptr<i32>>
+}
+
+// -----
+
+func.func @zero_vector() {
+ // expected-error @+1 {{the number of vector elements must be positive}}
+ "some.op"() : () -> !llvm.vec<0 x ptr<i32>>
+}
diff --git a/mlir/test/Dialect/LLVMIR/types-invalid.mlir b/mlir/test/Dialect/LLVMIR/types-invalid.mlir
index d8ac523b86d9..fce100e6a865 100644
--- a/mlir/test/Dialect/LLVMIR/types-invalid.mlir
+++ b/mlir/test/Dialect/LLVMIR/types-invalid.mlir
@@ -21,15 +21,8 @@ func.func @function_taking_function() {
// -----
-func.func @void_pointer() {
- // expected-error @+1 {{invalid pointer element type}}
- "some.op"() : () -> !llvm.ptr<void>
-}
-
-// -----
-
func.func @repeated_struct_name() {
- "some.op"() : () -> !llvm.struct<"a", (ptr<struct<"a">>)>
+ "some.op"() : () -> !llvm.struct<"a", (ptr)>
// expected-error @+1 {{identified type already used with a
diff erent body}}
"some.op"() : () -> !llvm.struct<"a", (i32)>
}
@@ -113,28 +106,28 @@ func.func @identified_struct_with_void() {
func.func @dynamic_vector() {
// expected-error @+1 {{expected '? x <integer> x <type>' or '<integer> x <type>'}}
- "some.op"() : () -> !llvm.vec<? x ptr<f32>>
+ "some.op"() : () -> !llvm.vec<? x ptr>
}
// -----
func.func @dynamic_scalable_vector() {
// expected-error @+1 {{expected '? x <integer> x <type>' or '<integer> x <type>'}}
- "some.op"() : () -> !llvm.vec<?x? x ptr<f32>>
+ "some.op"() : () -> !llvm.vec<?x? x ptr>
}
// -----
func.func @unscalable_vector() {
// expected-error @+1 {{expected '? x <integer> x <type>' or '<integer> x <type>'}}
- "some.op"() : () -> !llvm.vec<4x4 x ptr<i32>>
+ "some.op"() : () -> !llvm.vec<4x4 x ptr>
}
// -----
func.func @zero_vector() {
// expected-error @+1 {{the number of vector elements must be positive}}
- "some.op"() : () -> !llvm.vec<0 x ptr<i32>>
+ "some.op"() : () -> !llvm.vec<0 x ptr>
}
// -----
diff --git a/mlir/test/Dialect/LLVMIR/types-typed-pointers.mlir b/mlir/test/Dialect/LLVMIR/types-typed-pointers.mlir
new file mode 100644
index 000000000000..2d63f379c2ee
--- /dev/null
+++ b/mlir/test/Dialect/LLVMIR/types-typed-pointers.mlir
@@ -0,0 +1,118 @@
+// RUN: mlir-opt -allow-unregistered-dialect %s -split-input-file | mlir-opt -allow-unregistered-dialect | FileCheck %s
+
+// CHECK-LABEL: @ptr
+func.func @ptr() {
+ // CHECK: !llvm.ptr<i8>
+ "some.op"() : () -> !llvm.ptr<i8>
+ // CHECK: !llvm.ptr<f32>
+ "some.op"() : () -> !llvm.ptr<f32>
+ // CHECK: !llvm.ptr<ptr<i8>>
+ "some.op"() : () -> !llvm.ptr<ptr<i8>>
+ // CHECK: !llvm.ptr<ptr<ptr<ptr<ptr<i8>>>>>
+ "some.op"() : () -> !llvm.ptr<ptr<ptr<ptr<ptr<i8>>>>>
+ // CHECK: !llvm.ptr<i8>
+ "some.op"() : () -> !llvm.ptr<i8, 0>
+ // CHECK: !llvm.ptr<i8, 1>
+ "some.op"() : () -> !llvm.ptr<i8, 1>
+ // CHECK: !llvm.ptr<i8, 42>
+ "some.op"() : () -> !llvm.ptr<i8, 42>
+ // CHECK: !llvm.ptr<ptr<i8, 42>, 9>
+ "some.op"() : () -> !llvm.ptr<ptr<i8, 42>, 9>
+ // CHECK: !llvm.ptr
+ "some.op"() : () -> !llvm.ptr
+ // CHECK: !llvm.ptr<42>
+ "some.op"() : () -> !llvm.ptr<42>
+ return
+}
+
+// CHECK-LABEL: @vec
+func.func @vec() {
+ // CHECK: vector<4xi32>
+ "some.op"() : () -> vector<4xi32>
+ // CHECK: vector<4xf32>
+ "some.op"() : () -> vector<4xf32>
+ // CHECK: !llvm.vec<? x 4 x i32>
+ "some.op"() : () -> !llvm.vec<? x 4 x i32>
+ // CHECK: !llvm.vec<? x 8 x f16>
+ "some.op"() : () -> !llvm.vec<? x 8 x f16>
+ // CHECK: !llvm.vec<4 x ptr<i8>>
+ "some.op"() : () -> !llvm.vec<4 x ptr<i8>>
+ return
+}
+
+// CHECK-LABEL: @array
+func.func @array() {
+ // CHECK: !llvm.array<10 x i32>
+ "some.op"() : () -> !llvm.array<10 x i32>
+ // CHECK: !llvm.array<8 x f32>
+ "some.op"() : () -> !llvm.array<8 x f32>
+ // CHECK: !llvm.array<10 x ptr<i32, 4>>
+ "some.op"() : () -> !llvm.array<10 x ptr<i32, 4>>
+ // CHECK: !llvm.array<10 x array<4 x f32>>
+ "some.op"() : () -> !llvm.array<10 x array<4 x f32>>
+ return
+}
+
+// CHECK-LABEL: @identified_struct
+func.func @identified_struct() {
+ // CHECK: !llvm.struct<"empty", ()>
+ "some.op"() : () -> !llvm.struct<"empty", ()>
+ // CHECK: !llvm.struct<"opaque", opaque>
+ "some.op"() : () -> !llvm.struct<"opaque", opaque>
+ // CHECK: !llvm.struct<"long", (i32, struct<(i32, i1)>, f32, ptr<func<void ()>>)>
+ "some.op"() : () -> !llvm.struct<"long", (i32, struct<(i32, i1)>, f32, ptr<func<void ()>>)>
+ // CHECK: !llvm.struct<"self-recursive", (ptr<struct<"self-recursive">>)>
+ "some.op"() : () -> !llvm.struct<"self-recursive", (ptr<struct<"self-recursive">>)>
+ // CHECK: !llvm.struct<"unpacked", (i32)>
+ "some.op"() : () -> !llvm.struct<"unpacked", (i32)>
+ // CHECK: !llvm.struct<"packed", packed (i32)>
+ "some.op"() : () -> !llvm.struct<"packed", packed (i32)>
+ // CHECK: !llvm.struct<"name with spaces and !^$@$#", packed (i32)>
+ "some.op"() : () -> !llvm.struct<"name with spaces and !^$@$#", packed (i32)>
+
+ // CHECK: !llvm.struct<"mutually-a", (ptr<struct<"mutually-b", (ptr<struct<"mutually-a">, 3>)>>)>
+ "some.op"() : () -> !llvm.struct<"mutually-a", (ptr<struct<"mutually-b", (ptr<struct<"mutually-a">, 3>)>>)>
+ // CHECK: !llvm.struct<"mutually-b", (ptr<struct<"mutually-a", (ptr<struct<"mutually-b">>)>, 3>)>
+ "some.op"() : () -> !llvm.struct<"mutually-b", (ptr<struct<"mutually-a", (ptr<struct<"mutually-b">>)>, 3>)>
+ // CHECK: !llvm.struct<"referring-another", (ptr<struct<"unpacked", (i32)>>)>
+ "some.op"() : () -> !llvm.struct<"referring-another", (ptr<struct<"unpacked", (i32)>>)>
+
+ // CHECK: !llvm.struct<"struct-of-arrays", (array<10 x i32>)>
+ "some.op"() : () -> !llvm.struct<"struct-of-arrays", (array<10 x i32>)>
+ // CHECK: !llvm.array<10 x struct<"array-of-structs", (i32)>>
+ "some.op"() : () -> !llvm.array<10 x struct<"array-of-structs", (i32)>>
+ // CHECK: !llvm.ptr<struct<"ptr-to-struct", (i8)>>
+ "some.op"() : () -> !llvm.ptr<struct<"ptr-to-struct", (i8)>>
+ return
+}
+
+// CHECK-LABEL: @ptr_elem_interface
+// CHECK-COUNT-3: !llvm.ptr<!test.smpla>
+// CHECK: llvm.mlir.undef : !llvm.ptr<!test.smpla>
+func.func @ptr_elem_interface(%arg0: !llvm.ptr<!test.smpla>) {
+ %0 = llvm.load %arg0 : !llvm.ptr<!test.smpla>
+ llvm.store %0, %arg0 : !llvm.ptr<!test.smpla>
+ llvm.mlir.undef : !llvm.ptr<!test.smpla>
+ return
+}
+
+// -----
+
+// Check that type aliases can be used inside LLVM dialect types. Note that
+// currently they are _not_ printed back as this would require
+// DialectAsmPrinter to have a mechanism for querying the presence and
+// usability of an alias outside of its `printType` method.
+
+!baz = i64
+!qux = !llvm.struct<(!baz)>
+
+!rec = !llvm.struct<"a", (ptr<struct<"a">>)>
+
+// CHECK: aliases
+llvm.func @aliases() {
+ // CHECK: !llvm.struct<(i32, f32, struct<(i64)>)>
+ "some.op"() : () -> !llvm.struct<(i32, f32, !qux)>
+ // CHECK: !llvm.struct<"a", (ptr<struct<"a">>)>
+ "some.op"() : () -> !rec
+ llvm.return
+}
diff --git a/mlir/test/Dialect/LLVMIR/types.mlir b/mlir/test/Dialect/LLVMIR/types.mlir
index 54c44a6aa58a..42352ce697f0 100644
--- a/mlir/test/Dialect/LLVMIR/types.mlir
+++ b/mlir/test/Dialect/LLVMIR/types.mlir
@@ -57,26 +57,14 @@ func.func @integer() {
// CHECK-LABEL: @ptr
func.func @ptr() {
- // CHECK: !llvm.ptr<i8>
- "some.op"() : () -> !llvm.ptr<i8>
- // CHECK: !llvm.ptr<f32>
- "some.op"() : () -> !llvm.ptr<f32>
- // CHECK: !llvm.ptr<ptr<i8>>
- "some.op"() : () -> !llvm.ptr<ptr<i8>>
- // CHECK: !llvm.ptr<ptr<ptr<ptr<ptr<i8>>>>>
- "some.op"() : () -> !llvm.ptr<ptr<ptr<ptr<ptr<i8>>>>>
- // CHECK: !llvm.ptr<i8>
- "some.op"() : () -> !llvm.ptr<i8, 0>
- // CHECK: !llvm.ptr<i8, 1>
- "some.op"() : () -> !llvm.ptr<i8, 1>
- // CHECK: !llvm.ptr<i8, 42>
- "some.op"() : () -> !llvm.ptr<i8, 42>
- // CHECK: !llvm.ptr<ptr<i8, 42>, 9>
- "some.op"() : () -> !llvm.ptr<ptr<i8, 42>, 9>
// CHECK: !llvm.ptr
"some.op"() : () -> !llvm.ptr
+ // CHECK: !llvm.ptr
+ "some.op"() : () -> !llvm.ptr<0>
// CHECK: !llvm.ptr<42>
"some.op"() : () -> !llvm.ptr<42>
+ // CHECK: !llvm.ptr<ptr<42>, 9>
+ "some.op"() : () -> !llvm.ptr<ptr<42>, 9>
return
}
@@ -90,8 +78,8 @@ func.func @vec() {
"some.op"() : () -> !llvm.vec<? x 4 x i32>
// CHECK: !llvm.vec<? x 8 x f16>
"some.op"() : () -> !llvm.vec<? x 8 x f16>
- // CHECK: !llvm.vec<4 x ptr<i8>>
- "some.op"() : () -> !llvm.vec<4 x ptr<i8>>
+ // CHECK: !llvm.vec<4 x ptr>
+ "some.op"() : () -> !llvm.vec<4 x ptr>
return
}
@@ -101,8 +89,8 @@ func.func @array() {
"some.op"() : () -> !llvm.array<10 x i32>
// CHECK: !llvm.array<8 x f32>
"some.op"() : () -> !llvm.array<8 x f32>
- // CHECK: !llvm.array<10 x ptr<i32, 4>>
- "some.op"() : () -> !llvm.array<10 x ptr<i32, 4>>
+ // CHECK: !llvm.array<10 x ptr<4>>
+ "some.op"() : () -> !llvm.array<10 x ptr<4>>
// CHECK: !llvm.array<10 x array<4 x f32>>
"some.op"() : () -> !llvm.array<10 x array<4 x f32>>
return
@@ -147,30 +135,22 @@ func.func @identified_struct() {
"some.op"() : () -> !llvm.struct<"empty", ()>
// CHECK: !llvm.struct<"opaque", opaque>
"some.op"() : () -> !llvm.struct<"opaque", opaque>
- // CHECK: !llvm.struct<"long", (i32, struct<(i32, i1)>, f32, ptr<func<void ()>>)>
- "some.op"() : () -> !llvm.struct<"long", (i32, struct<(i32, i1)>, f32, ptr<func<void ()>>)>
- // CHECK: !llvm.struct<"self-recursive", (ptr<struct<"self-recursive">>)>
- "some.op"() : () -> !llvm.struct<"self-recursive", (ptr<struct<"self-recursive">>)>
+ // CHECK: !llvm.struct<"long", (i32, struct<(i32, i1)>, f32, ptr)>
+ "some.op"() : () -> !llvm.struct<"long", (i32, struct<(i32, i1)>, f32, ptr)>
// CHECK: !llvm.struct<"unpacked", (i32)>
"some.op"() : () -> !llvm.struct<"unpacked", (i32)>
// CHECK: !llvm.struct<"packed", packed (i32)>
"some.op"() : () -> !llvm.struct<"packed", packed (i32)>
// CHECK: !llvm.struct<"name with spaces and !^$@$#", packed (i32)>
"some.op"() : () -> !llvm.struct<"name with spaces and !^$@$#", packed (i32)>
-
- // CHECK: !llvm.struct<"mutually-a", (ptr<struct<"mutually-b", (ptr<struct<"mutually-a">, 3>)>>)>
- "some.op"() : () -> !llvm.struct<"mutually-a", (ptr<struct<"mutually-b", (ptr<struct<"mutually-a">, 3>)>>)>
- // CHECK: !llvm.struct<"mutually-b", (ptr<struct<"mutually-a", (ptr<struct<"mutually-b">>)>, 3>)>
- "some.op"() : () -> !llvm.struct<"mutually-b", (ptr<struct<"mutually-a", (ptr<struct<"mutually-b">>)>, 3>)>
- // CHECK: !llvm.struct<"referring-another", (ptr<struct<"unpacked", (i32)>>)>
- "some.op"() : () -> !llvm.struct<"referring-another", (ptr<struct<"unpacked", (i32)>>)>
-
+ // CHECK: !llvm.struct<"outer", (struct<"nested", ()>)>
+ "some.op"() : () -> !llvm.struct<"outer", (struct<"nested", ()>)>
+ // CHECK: !llvm.struct<"referring-another", (ptr)>
+ "some.op"() : () -> !llvm.struct<"referring-another", (ptr)>
// CHECK: !llvm.struct<"struct-of-arrays", (array<10 x i32>)>
"some.op"() : () -> !llvm.struct<"struct-of-arrays", (array<10 x i32>)>
// CHECK: !llvm.array<10 x struct<"array-of-structs", (i32)>>
"some.op"() : () -> !llvm.array<10 x struct<"array-of-structs", (i32)>>
- // CHECK: !llvm.ptr<struct<"ptr-to-struct", (i8)>>
- "some.op"() : () -> !llvm.ptr<struct<"ptr-to-struct", (i8)>>
return
}
@@ -180,16 +160,6 @@ func.func @verbose() {
return
}
-// CHECK-LABEL: @ptr_elem_interface
-// CHECK-COUNT-3: !llvm.ptr<!test.smpla>
-// CHECK: llvm.mlir.undef : !llvm.ptr<!test.smpla>
-func.func @ptr_elem_interface(%arg0: !llvm.ptr<!test.smpla>) {
- %0 = llvm.load %arg0 : !llvm.ptr<!test.smpla>
- llvm.store %0, %arg0 : !llvm.ptr<!test.smpla>
- llvm.mlir.undef : !llvm.ptr<!test.smpla>
- return
-}
-
// -----
// Check that type aliases can be used inside LLVM dialect types. Note that
@@ -200,13 +170,9 @@ func.func @ptr_elem_interface(%arg0: !llvm.ptr<!test.smpla>) {
!baz = i64
!qux = !llvm.struct<(!baz)>
-!rec = !llvm.struct<"a", (ptr<struct<"a">>)>
-
// CHECK: aliases
llvm.func @aliases() {
// CHECK: !llvm.struct<(i32, f32, struct<(i64)>)>
"some.op"() : () -> !llvm.struct<(i32, f32, !qux)>
- // CHECK: !llvm.struct<"a", (ptr<struct<"a">>)>
- "some.op"() : () -> !rec
llvm.return
}
More information about the Mlir-commits
mailing list