diff --git a/mlir/test/Dialect/LLVMIR/callgraph.mlir b/mlir/test/Dialect/LLVMIR/callgraph.mlir --- a/mlir/test/Dialect/LLVMIR/callgraph.mlir +++ b/mlir/test/Dialect/LLVMIR/callgraph.mlir @@ -58,33 +58,32 @@ // CHECK-DAG: -- Call-Edge : // CHECK: -- SCCs -- - llvm.mlir.global external constant @_ZTIi() : !llvm.ptr + llvm.mlir.global external constant @_ZTIi() : !llvm.ptr llvm.func @foo(%arg0: i32) -> !llvm.struct<(i32, f64, i32)> - llvm.func @bar(!llvm.ptr, !llvm.ptr, !llvm.ptr) + 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> - %4 = llvm.mlir.null : !llvm.ptr - %5 = llvm.mlir.addressof @_ZTIi : !llvm.ptr> - %6 = llvm.bitcast %5 : !llvm.ptr> to !llvm.ptr - %7 = llvm.mlir.constant(1 : i32) : i32 - %8 = llvm.alloca %7 x i8 : (i32) -> !llvm.ptr - %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>) (catch %6 : !llvm.ptr) (filter %2 : !llvm.array<1 x i8>) : !llvm.struct<(ptr, i32)> - %11 = llvm.intr.eh.typeid.for %6 : (!llvm.ptr) -> i32 - llvm.resume %10 : !llvm.struct<(ptr, 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, !llvm.ptr, !llvm.ptr) -> () + 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.mlir b/mlir/test/Dialect/LLVMIR/canonicalize-typed-pointers.mlir copy from mlir/test/Dialect/LLVMIR/canonicalize.mlir copy to mlir/test/Dialect/LLVMIR/canonicalize-typed-pointers.mlir --- a/mlir/test/Dialect/LLVMIR/canonicalize.mlir +++ b/mlir/test/Dialect/LLVMIR/canonicalize-typed-pointers.mlir @@ -1,58 +1,5 @@ // RUN: mlir-opt --pass-pipeline='builtin.module(llvm.func(canonicalize{test-convergence}))' %s -split-input-file | FileCheck %s -// CHECK-LABEL: fold_extractvalue -llvm.func @fold_extractvalue() -> i32 { - // CHECK-DAG: %[[C0:.*]] = arith.constant 0 : i32 - %c0 = arith.constant 0 : i32 - // CHECK-DAG: %[[C1:.*]] = arith.constant 1 : i32 - %c1 = arith.constant 1 : i32 - - %0 = llvm.mlir.undef : !llvm.struct<(i32, i32)> - - // CHECK-NOT: insertvalue - %1 = llvm.insertvalue %c0, %0[0] : !llvm.struct<(i32, i32)> - %2 = llvm.insertvalue %c1, %1[1] : !llvm.struct<(i32, i32)> - - // CHECK-NOT: extractvalue - %3 = llvm.extractvalue %2[0] : !llvm.struct<(i32, i32)> - %4 = llvm.extractvalue %2[1] : !llvm.struct<(i32, i32)> - - // CHECK: llvm.add %[[C0]], %[[C1]] - %5 = llvm.add %3, %4 : i32 - llvm.return %5 : i32 -} - -// ----- - -// CHECK-LABEL: no_fold_extractvalue -llvm.func @no_fold_extractvalue(%arr: !llvm.array<4 x f32>) -> f32 { - %f0 = arith.constant 0.0 : f32 - %0 = llvm.mlir.undef : !llvm.array<4 x !llvm.array<4 x f32>> - - // CHECK: insertvalue - // CHECK: insertvalue - // CHECK: extractvalue - %1 = llvm.insertvalue %f0, %0[0, 0] : !llvm.array<4 x !llvm.array<4 x f32>> - %2 = llvm.insertvalue %arr, %1[0] : !llvm.array<4 x !llvm.array<4 x 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 -llvm.func @fold_unrelated_extractvalue(%arr: !llvm.array<4 x f32>) -> f32 { - %f0 = arith.constant 0.0 : f32 - // CHECK-NOT: insertvalue - // CHECK: extractvalue - %2 = llvm.insertvalue %f0, %arr[0] : !llvm.array<4 x f32> - %3 = llvm.extractvalue %2[1] : !llvm.array<4 x f32> - llvm.return %3 : f32 -} - -// ----- - // CHECK-LABEL: fold_bitcast // CHECK-SAME: %[[a0:arg[0-9]+]] // CHECK-NEXT: llvm.return %[[a0]] @@ -100,15 +47,7 @@ llvm.return %c : !llvm.ptr } -// CHECK-LABEL: fold_gep_neg -// CHECK-SAME: %[[a0:arg[0-9]+]] -// CHECK-NEXT: %[[RES:.*]] = llvm.getelementptr inbounds %[[a0]][0, 1] -// CHECK-NEXT: llvm.return %[[RES]] -llvm.func @fold_gep_neg(%x : !llvm.ptr) -> !llvm.ptr { - %c0 = arith.constant 0 : i32 - %0 = llvm.getelementptr inbounds %x[%c0, 1] : (!llvm.ptr, i32) -> !llvm.ptr, !llvm.struct<(i32, i32)> - llvm.return %0 : !llvm.ptr -} +// ----- // CHECK-LABEL: fold_gep_canon // CHECK-SAME: %[[a0:arg[0-9]+]] @@ -120,24 +59,6 @@ llvm.return %c : !llvm.ptr } - -// ----- - -// Check that LLVM constants participate in cross-dialect constant folding. The -// resulting constant is created in the arith dialect because the last folded -// operation belongs to it. -// CHECK-LABEL: llvm_constant -llvm.func @llvm_constant() -> i32 { - // CHECK-NOT: llvm.mlir.constant - %0 = llvm.mlir.constant(40 : i32) : i32 - %1 = llvm.mlir.constant(42 : i32) : i32 - // CHECK: %[[RES:.*]] = arith.constant 82 : i32 - // CHECK-NOT: arith.addi - %2 = arith.addi %0, %1 : i32 - // CHECK: return %[[RES]] - llvm.return %2 : i32 -} - // ----- // CHECK-LABEL: load_dce diff --git a/mlir/test/Dialect/LLVMIR/canonicalize.mlir b/mlir/test/Dialect/LLVMIR/canonicalize.mlir --- a/mlir/test/Dialect/LLVMIR/canonicalize.mlir +++ b/mlir/test/Dialect/LLVMIR/canonicalize.mlir @@ -37,8 +37,8 @@ %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 @@ // CHECK-LABEL: fold_bitcast // CHECK-SAME: %[[a0:arg[0-9]+]] // CHECK-NEXT: llvm.return %[[a0]] -llvm.func @fold_bitcast(%x : !llvm.ptr) -> !llvm.ptr { - %c = llvm.bitcast %x : !llvm.ptr to !llvm.ptr - llvm.return %c : !llvm.ptr +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) -> !llvm.ptr { - %c = llvm.bitcast %x : !llvm.ptr to !llvm.ptr - %d = llvm.bitcast %c : !llvm.ptr to !llvm.ptr - llvm.return %d : !llvm.ptr +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 @@ // CHECK-LABEL: fold_addrcast // CHECK-SAME: %[[a0:arg[0-9]+]] // CHECK-NEXT: llvm.return %[[a0]] -llvm.func @fold_addrcast(%x : !llvm.ptr) -> !llvm.ptr { - %c = llvm.addrspacecast %x : !llvm.ptr to !llvm.ptr - llvm.return %c : !llvm.ptr +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) -> !llvm.ptr { - %c = llvm.addrspacecast %x : !llvm.ptr to !llvm.ptr - %d = llvm.addrspacecast %c : !llvm.ptr to !llvm.ptr - llvm.return %d : !llvm.ptr +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 @@ // CHECK-LABEL: fold_gep // CHECK-SAME: %[[a0:arg[0-9]+]] // CHECK-NEXT: llvm.return %[[a0]] -llvm.func @fold_gep(%x : !llvm.ptr) -> !llvm.ptr { +llvm.func @fold_gep(%x : !llvm.ptr) -> !llvm.ptr { %c0 = arith.constant 0 : i32 - %c = llvm.getelementptr %x[%c0] : (!llvm.ptr, i32) -> !llvm.ptr - llvm.return %c : !llvm.ptr + %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 @@ // 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) -> !llvm.ptr { +llvm.func @fold_gep_canon(%x : !llvm.ptr) -> !llvm.ptr { %c2 = arith.constant 2 : i32 - %c = llvm.getelementptr %x[%c2] : (!llvm.ptr, i32) -> !llvm.ptr - llvm.return %c : !llvm.ptr + %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 @@ // CHECK-LABEL: load_dce // CHECK-NEXT: llvm.return -llvm.func @load_dce(%x : !llvm.ptr) { - %0 = llvm.load %x : !llvm.ptr +llvm.func @load_dce(%x : !llvm.ptr) { + %0 = llvm.load %x : !llvm.ptr -> i8 llvm.return } -llvm.mlir.global external @fp() : !llvm.ptr +llvm.mlir.global external @fp() : !llvm.ptr // CHECK-LABEL: addr_dce // CHECK-NEXT: llvm.return -llvm.func @addr_dce(%x : !llvm.ptr) { - %0 = llvm.mlir.addressof @fp : !llvm.ptr> +llvm.func @addr_dce(%x : !llvm.ptr) { + %0 = llvm.mlir.addressof @fp : !llvm.ptr llvm.return } @@ -160,6 +159,6 @@ // CHECK-NEXT: llvm.return llvm.func @alloca_dce() { %c1_i64 = arith.constant 1 : i64 - %0 = llvm.alloca %c1_i64 x i32 : (i64) -> !llvm.ptr + %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 --- 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 + %alloc = llvm.alloca %allocCount x i64 : (i32) -> !llvm.ptr // CHECK: llvm.intr.dbg.declare #[[VAR0]] = %[[ALLOC]] - llvm.intr.dbg.declare #var0 = %alloc : !llvm.ptr + llvm.intr.dbg.declare #var0 = %alloc : !llvm.ptr llvm.return } diff --git a/mlir/test/Dialect/LLVMIR/dynamic-gep-index.mlir b/mlir/test/Dialect/LLVMIR/dynamic-gep-index-typed-pointers.mlir copy from mlir/test/Dialect/LLVMIR/dynamic-gep-index.mlir copy to mlir/test/Dialect/LLVMIR/dynamic-gep-index-typed-pointers.mlir diff --git a/mlir/test/Dialect/LLVMIR/dynamic-gep-index.mlir b/mlir/test/Dialect/LLVMIR/dynamic-gep-index.mlir --- 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 : vector<2xi32>>, #dlti.dl_entry : vector<2xi32>>, #dlti.dl_entry : vector<2xi32>>, #dlti.dl_entry : vector<2xi32>>, #dlti.dl_entry : vector<2xi32>>, #dlti.dl_entry : vector<2xi32>>, #dlti.dl_entry : vector<2xi32>>, #dlti.dl_entry : vector<2xi32>>>} { - // CHECK: llvm.func @foo(%[[ARG0:.+]]: !llvm.ptr>, %[[ARG1:.+]]: i32) - llvm.func @foo(%arg0: !llvm.ptr, 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} : (!llvm.ptr, array<4 x i32>)>>, i32, i32) -> !llvm.ptr + %1 = "llvm.getelementptr"(%arg0, %0, %arg1) {elem_type = !llvm.struct<"my_struct", (struct<"sub_struct", (i32, i8)>, array<4 x i32>)>, rawConstantIndices = array} : (!llvm.ptr, i32, i32) -> !llvm.ptr llvm.return } } diff --git a/mlir/test/Dialect/LLVMIR/func.mlir b/mlir/test/Dialect/LLVMIR/func.mlir --- a/mlir/test/Dialect/LLVMIR/func.mlir +++ b/mlir/test/Dialect/LLVMIR/func.mlir @@ -33,10 +33,10 @@ // GENERIC-SAME: () -> () }) {sym_name = "baz", function_type = !llvm.func} : () -> () - // CHECK: llvm.func @qux(!llvm.ptr {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, i64)>, + }) {sym_name = "qux", function_type = !llvm.func, arg_attrs = [{llvm.noalias}, {}], xxx = {yyy = 42}} : () -> () // CHECK: llvm.func @roundtrip1() @@ -71,56 +71,56 @@ // CHECK: llvm.func @roundtrip8() -> i32 llvm.func @roundtrip8() -> i32 attributes {} - // CHECK: llvm.func @roundtrip9(!llvm.ptr {llvm.noalias}) - llvm.func @roundtrip9(!llvm.ptr {llvm.noalias}) + // CHECK: llvm.func @roundtrip9(!llvm.ptr {llvm.noalias}) + llvm.func @roundtrip9(!llvm.ptr {llvm.noalias}) - // CHECK: llvm.func @roundtrip10(!llvm.ptr {llvm.noalias}) - llvm.func @roundtrip10(%arg0: !llvm.ptr {llvm.noalias}) + // CHECK: llvm.func @roundtrip10(!llvm.ptr {llvm.noalias}) + llvm.func @roundtrip10(%arg0: !llvm.ptr {llvm.noalias}) - // CHECK: llvm.func @roundtrip11(%{{.*}}: !llvm.ptr {llvm.noalias}) { - llvm.func @roundtrip11(%arg0: !llvm.ptr {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 {llvm.noalias}) + // CHECK: llvm.func @roundtrip12(%{{.*}}: !llvm.ptr {llvm.noalias}) // CHECK: attributes {foo = 42 : i32} - llvm.func @roundtrip12(%arg0: !llvm.ptr {llvm.noalias}) + llvm.func @roundtrip12(%arg0: !llvm.ptr {llvm.noalias}) attributes {foo = 42 : i32} { llvm.return } - // CHECK: llvm.func @byvalattr(%{{.*}}: !llvm.ptr {llvm.byval = i32}) - llvm.func @byvalattr(%arg0: !llvm.ptr {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 {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")) { + // 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 {llvm.nest}) - llvm.func @nestattr(%arg0: !llvm.ptr {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 {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 @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 {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 @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 --- /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 + %0 = llvm.mlir.addressof @global : !llvm.ptr + + // CHECK: llvm.mlir.addressof @".string" : !llvm.ptr> + %1 = llvm.mlir.addressof @".string" : !llvm.ptr> + + 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 + 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 + 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 + llvm.return +} diff --git a/mlir/test/Dialect/LLVMIR/global.mlir b/mlir/test/Dialect/LLVMIR/global.mlir --- a/mlir/test/Dialect/LLVMIR/global.mlir +++ b/mlir/test/Dialect/LLVMIR/global.mlir @@ -66,17 +66,14 @@ // CHECK-LABEL: references func.func @references() { - // CHECK: llvm.mlir.addressof @global : !llvm.ptr - %0 = llvm.mlir.addressof @global : !llvm.ptr - - // CHECK: llvm.mlir.addressof @".string" : !llvm.ptr> - %1 = llvm.mlir.addressof @".string" : !llvm.ptr> + // 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 @@ // 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> + llvm.mlir.addressof "foo" : i64 : !llvm.ptr llvm.return } @@ -172,27 +169,7 @@ func.func @foo() { // expected-error @+1 {{must reference a global defined by 'llvm.mlir.global'}} - llvm.mlir.addressof @foo : !llvm.ptr> - 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 - 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 + llvm.mlir.addressof @foo : !llvm.ptr llvm.return } @@ -224,23 +201,15 @@ 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 + 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 - 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 --- /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@below {{unexpected 'elem_type' attribute when non-opaque pointer type is used}} + "llvm.alloca"(%sz) { elem_type = i32 } : (i64) -> !llvm.ptr +} + +// ----- + +func.func @gep_missing_input_type(%pos : i64, %base : !llvm.ptr) { + // expected-error@+1 {{2 operands present, but expected 0}} + llvm.getelementptr %base[%pos] : () -> (!llvm.ptr) +} + +// ----- + +func.func @gep_missing_result_type(%pos : i64, %base : !llvm.ptr) { + // expected-error@+1 {{op requires one result}} + llvm.getelementptr %base[%pos] : (!llvm.ptr, i64) -> () +} + +// ----- + +func.func @gep_non_function_type(%pos : i64, %base : !llvm.ptr) { + // expected-error@+1 {{invalid kind of type specified}} + llvm.getelementptr %base[%pos] : !llvm.ptr +} + +// ----- + +func.func @gep_too_few_dynamic(%base : !llvm.ptr) { + // expected-error@+1 {{expected as many dynamic indices as specified in 'rawConstantIndices'}} + %1 = "llvm.getelementptr"(%base) {rawConstantIndices = array} : (!llvm.ptr) -> !llvm.ptr +} + +// ----- + +func.func @call_variadic(%callee : !llvm.ptr>, %arg : i8) { + // expected-error@+1 {{indirect calls to variadic functions are not supported}} + llvm.call %callee(%arg) : !llvm.ptr>, (i8) -> (i8) + llvm.return +} + +// ----- + +func.func @indirect_callee_arg_mismatch(%arg0 : i32, %callee : !llvm.ptr>) { + // expected-error@+1 {{'llvm.call' op operand type mismatch for operand 0: 'i32' != 'i8'}} + "llvm.call"(%callee, %arg0) : (!llvm.ptr>, i32) -> () + llvm.return +} + +// ----- + +func.func @indirect_callee_return_mismatch(%callee : !llvm.ptr>) { + // expected-error@+1 {{'llvm.call' op result type mismatch: 'i32' != 'i8'}} + "llvm.call"(%callee) : (!llvm.ptr>) -> (i32) + llvm.return +} + +// ----- + +func.func @atomicrmw_mismatched_operands(%f32_ptr : !llvm.ptr, %i32 : i32) { + // expected-error@+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, i32) -> i32 + llvm.return +} + +// ----- + +func.func @cmpxchg_expected_ptr(%f32 : f32) { + // expected-error@+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, %i32 : i32) { + // expected-error@+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, 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>) -> 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@+1 {{clause #0 is not a known constant - null, addressof, bitcast}} + %3 = llvm.landingpad cleanup (catch %1 : i32) (catch %arg0 : !llvm.ptr>) : !llvm.struct<(ptr, 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 : (i32) -> !llvm.ptr> + // expected-note@+1 {{global addresses expected as operand to bitcast used in clauses for landingpad}} + %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@+1 {{constant clauses expected}} + %5 = llvm.landingpad (catch %2 : !llvm.ptr) : !llvm.struct<(ptr, 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@+1 {{landingpad instruction expects at least one clause or cleanup attribute}} + %2 = llvm.landingpad : !llvm.struct<(ptr, 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, i32)> + // expected-error@+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@+1 {{llvm.landingpad needs to be in a function with a personality}} + %2 = llvm.landingpad cleanup : !llvm.struct<(ptr, i32)> + llvm.resume %2 : !llvm.struct<(ptr, i32)> +} + +// ----- + +llvm.func @wmmaLoadOp_invalid_mem_space(%arg0: !llvm.ptr<5>, %arg1: i32) { + // expected-error@+1 {{'nvvm.wmma.load' op expected source pointer in memory space 0, 1, 3}} + %0 = nvvm.wmma.load %arg0, %arg1 + {eltype = #nvvm.mma_type, frag = #nvvm.mma_frag, k = 16 : i32, layout = #nvvm.mma_layout, 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@+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, frag = #nvvm.mma_frag, k = 16 : i32, layout = #nvvm.mma_layout, 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@+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, frag = #nvvm.mma_frag, k = 16 : i32, layout = #nvvm.mma_layout, 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@+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, frag = #nvvm.mma_frag, k = 16 : i32, layout = #nvvm.mma_layout, 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@+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, k = 16 : i32, layout = #nvvm.mma_layout, 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) { + // expected-error@+1 {{'nvvm.ldmatrix' op expected source pointer in memory space 3}} + %l = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout} : (!llvm.ptr) -> i32 + llvm.return +} + +// ----- + +llvm.func @wmmald_matrix(%arg0: !llvm.ptr) { + // expected-error@+1 {{'nvvm.ldmatrix' op expected num attribute to be 1, 2 or 4}} + %l = nvvm.ldmatrix %arg0 {num = 3 : i32, layout = #nvvm.mma_layout} : (!llvm.ptr) -> i32 + llvm.return +} + +// ----- + +llvm.func @wmmald_matrix(%arg0: !llvm.ptr) { + // expected-error@+1 {{'nvvm.ldmatrix' op expected destination type is i32}} + %l = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout} : (!llvm.ptr) -> !llvm.struct<(i32)> + llvm.return +} + +// ----- + +llvm.func @wmmald_matrix(%arg0: !llvm.ptr) { + // expected-error@+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} : (!llvm.ptr) -> !llvm.struct<(i32, i32)> + llvm.return +} + +// ----- + +func.func @cp_async(%arg0: !llvm.ptr, %arg1: !llvm.ptr) { + // expected-error @below {{expected byte size to be either 4, 8 or 16.}} + nvvm.cp.async.shared.global %arg0, %arg1, 32 : !llvm.ptr, !llvm.ptr + return +} + +// ----- + +func.func @cp_async(%arg0: !llvm.ptr, %arg1: !llvm.ptr) { + // expected-error @below {{bypass l1 is only support for 16 bytes copy.}} + nvvm.cp.async.shared.global %arg0, %arg1, 8 {bypass_l1} : !llvm.ptr, !llvm.ptr + return +} + +// ----- + +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>, i32, i32) -> !llvm.ptr + return +} + +// ----- + +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)>>, i64) -> !llvm.ptr + return +} diff --git a/mlir/test/Dialect/LLVMIR/invalid.mlir b/mlir/test/Dialect/LLVMIR/invalid.mlir --- a/mlir/test/Dialect/LLVMIR/invalid.mlir +++ b/mlir/test/Dialect/LLVMIR/invalid.mlir @@ -64,7 +64,7 @@ func.func @alloca_missing_input_type() { // expected-error@+1 {{expected trailing function type with one argument and one result}} - llvm.alloca %size x i32 : () -> (!llvm.ptr) + llvm.alloca %size x i32 : () -> (!llvm.ptr) } // ----- @@ -78,14 +78,14 @@ func.func @alloca_non_function_type() { // expected-error@+1 {{expected trailing function type with one argument and one result}} - llvm.alloca %size x i32 : !llvm.ptr + llvm.alloca %size x i32 : !llvm.ptr } // ----- func.func @alloca_non_integer_alignment() { // expected-error@+1 {{expected integer alignment}} - llvm.alloca %size x i32 {alignment = 3.0} : !llvm.ptr + llvm.alloca %size x i32 {alignment = 3.0} : !llvm.ptr } // ----- @@ -97,44 +97,37 @@ // ----- -func.func @alloca_ptr_type_attr_non_opaque_ptr(%sz : i64) { - // expected-error@below {{unexpected 'elem_type' attribute when non-opaque pointer type is used}} - "llvm.alloca"(%sz) { elem_type = i32 } : (i64) -> !llvm.ptr -} - -// ----- - -func.func @gep_missing_input_result_type(%pos : i64, %base : !llvm.ptr) { +func.func @gep_missing_input_result_type(%pos : i64, %base : !llvm.ptr) { // expected-error@+1 {{2 operands present, but expected 0}} llvm.getelementptr %base[%pos] : () -> () } // ----- -func.func @gep_missing_input_type(%pos : i64, %base : !llvm.ptr) { +func.func @gep_missing_input_type(%pos : i64, %base : !llvm.ptr) { // expected-error@+1 {{2 operands present, but expected 0}} - llvm.getelementptr %base[%pos] : () -> (!llvm.ptr) + llvm.getelementptr %base[%pos] : () -> (!llvm.ptr) } // ----- -func.func @gep_missing_result_type(%pos : i64, %base : !llvm.ptr) { +func.func @gep_missing_result_type(%pos : i64, %base : !llvm.ptr) { // expected-error@+1 {{op requires one result}} - llvm.getelementptr %base[%pos] : (!llvm.ptr, i64) -> () + llvm.getelementptr %base[%pos] : (!llvm.ptr, i64) -> () } // ----- -func.func @gep_non_function_type(%pos : i64, %base : !llvm.ptr) { +func.func @gep_non_function_type(%pos : i64, %base : !llvm.ptr) { // expected-error@+1 {{invalid kind of type specified}} - llvm.getelementptr %base[%pos] : !llvm.ptr + llvm.getelementptr %base[%pos] : !llvm.ptr } // ----- -func.func @gep_too_few_dynamic(%base : !llvm.ptr) { +func.func @gep_too_few_dynamic(%base : !llvm.ptr) { // expected-error@+1 {{expected as many dynamic indices as specified in 'rawConstantIndices'}} - %1 = "llvm.getelementptr"(%base) {rawConstantIndices = array} : (!llvm.ptr) -> !llvm.ptr + %1 = "llvm.getelementptr"(%base) {elem_type = f32, rawConstantIndices = array} : (!llvm.ptr) -> !llvm.ptr } // ----- @@ -302,14 +295,6 @@ // ----- -func.func @call_variadic(%callee : !llvm.ptr>, %arg : i8) { - // expected-error@+1 {{indirect calls to variadic functions are not supported}} - llvm.call %callee(%arg) : !llvm.ptr>, (i8) -> (i8) - llvm.return -} - -// ----- - func.func private @standard_func_callee() func.func @call_non_llvm() { @@ -346,14 +331,6 @@ // ----- -func.func @indirect_callee_arg_mismatch(%arg0 : i32, %callee : !llvm.ptr>) { - // expected-error@+1 {{'llvm.call' op operand type mismatch for operand 0: 'i32' != 'i8'}} - "llvm.call"(%callee, %arg0) : (!llvm.ptr>, i32) -> () - llvm.return -} - -// ----- - llvm.func @callee_func() -> (i8) func.func @callee_return_mismatch() { @@ -364,14 +341,6 @@ // ----- -func.func @indirect_callee_return_mismatch(%callee : !llvm.ptr>) { - // expected-error@+1 {{'llvm.call' op result type mismatch: 'i32' != 'i8'}} - "llvm.call"(%callee) : (!llvm.ptr>) -> (i32) - llvm.return -} - -// ----- - func.func @call_too_many_results(%callee : !llvm.ptr) { // expected-error@+1 {{expected function with 0 or 1 result}} llvm.call %callee() : !llvm.ptr, () -> (i32, i32) @@ -406,14 +375,14 @@ func.func @constant_wrong_type() { // expected-error@+1 {{only supports integer, float, string or elements attributes}} - llvm.mlir.constant(@constant_wrong_type) : !llvm.ptr> + llvm.mlir.constant(@constant_wrong_type) : !llvm.ptr } // ----- func.func @constant_wrong_type_string() { // expected-error@below {{expected array type of 3 i8 elements for the string constant}} - llvm.mlir.constant("foo") : !llvm.ptr + llvm.mlir.constant("foo") : !llvm.ptr } // ----- @@ -671,47 +640,39 @@ // ----- -func.func @atomicrmw_mismatched_operands(%f32_ptr : !llvm.ptr, %i32 : i32) { - // expected-error@+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, i32) -> i32 - llvm.return -} - -// ----- - -func.func @atomicrmw_mismatched_operands(%f32_ptr : !llvm.ptr, %f32 : f32) { +func.func @atomicrmw_mismatched_operands(%f32_ptr : !llvm.ptr, %f32 : f32) { // expected-error@+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) -> 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) { +func.func @atomicrmw_expected_float(%i32_ptr : !llvm.ptr, %i32 : i32) { // expected-error@+1 {{expected LLVM IR floating point type}} - %0 = llvm.atomicrmw fadd %i32_ptr, %i32 unordered : !llvm.ptr, 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) { +func.func @atomicrmw_unexpected_xchg_type(%i1_ptr : !llvm.ptr, %i1 : i1) { // expected-error@+1 {{unexpected LLVM IR type for 'xchg' bin_op}} - %0 = llvm.atomicrmw xchg %i1_ptr, %i1 unordered : !llvm.ptr, 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) { +func.func @atomicrmw_expected_int(%f32_ptr : !llvm.ptr, %f32 : f32) { // expected-error@+1 {{expected LLVM IR integer type}} - %0 = llvm.atomicrmw max %f32_ptr, %f32 unordered : !llvm.ptr, 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) { +func.func @cmpxchg_expected_ptr(%f32 : f32) { // expected-error@+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_mismatched_operands(%i64_ptr : !llvm.ptr, %i32 : i32) { - // expected-error@+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, i32, i32) -> !llvm.struct<(i32, i1)> - llvm.return -} - -// ----- - func.func @cmpxchg_mismatched_value_operands(%ptr : !llvm.ptr, %i32 : i32, %i64 : i64) { // expected-error@+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_unexpected_type(%i1_ptr : !llvm.ptr, %i1 : i1) { +func.func @cmpxchg_unexpected_type(%i1_ptr : !llvm.ptr, %i1 : i1) { // expected-error@+1 {{unexpected LLVM IR type}} - %0 = llvm.cmpxchg %i1_ptr, %i1, %i1 monotonic monotonic : !llvm.ptr, 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) { +func.func @cmpxchg_at_least_monotonic_success(%i32_ptr : !llvm.ptr, %i32 : i32) { // expected-error@+1 {{ordering must be at least 'monotonic'}} - %0 = llvm.cmpxchg %i32_ptr, %i32, %i32 unordered monotonic : !llvm.ptr, 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) { +func.func @cmpxchg_at_least_monotonic_failure(%i32_ptr : !llvm.ptr, %i32 : i32) { // expected-error@+1 {{ordering must be at least 'monotonic'}} - %0 = llvm.cmpxchg %i32_ptr, %i32, %i32 monotonic unordered : !llvm.ptr, 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) { +func.func @cmpxchg_failure_release(%i32_ptr : !llvm.ptr, %i32 : i32) { // expected-error@+1 {{failure ordering cannot be 'release' or 'acq_rel'}} - %0 = llvm.cmpxchg %i32_ptr, %i32, %i32 acq_rel release : !llvm.ptr, 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) { +func.func @cmpxchg_failure_acq_rel(%i32_ptr : !llvm.ptr, %i32 : i32) { // expected-error@+1 {{failure ordering cannot be 'release' or 'acq_rel'}} - %0 = llvm.cmpxchg %i32_ptr, %i32, %i32 acq_rel acq_rel : !llvm.ptr, i32 + %0 = llvm.cmpxchg %i32_ptr, %i32, %i32 acq_rel acq_rel : !llvm.ptr, i32 llvm.return } @@ -786,7 +739,7 @@ llvm.func @foo(i32) -> i32 llvm.func @__gxx_personality_v0(...) -> i32 -llvm.func @bad_landingpad(%arg0: !llvm.ptr>) -> 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.return %1 : i32 ^bb2: // pred: ^bb0 // expected-error@+1 {{clause #0 is not a known constant - null, addressof, bitcast}} - %3 = llvm.landingpad cleanup (catch %1 : i32) (catch %arg0 : !llvm.ptr>) : !llvm.struct<(ptr, 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 @caller(%arg0: i32) -> i32 attributes { personality = @__gxx_personality_v0} { %0 = llvm.mlir.constant(1 : i32) : i32 - %1 = llvm.alloca %0 x !llvm.ptr : (i32) -> !llvm.ptr> + %1 = llvm.alloca %0 x !llvm.ptr : (i32) -> !llvm.ptr // expected-note@+1 {{global addresses expected as operand to bitcast used in clauses for landingpad}} - %2 = llvm.bitcast %1 : !llvm.ptr> to !llvm.ptr + %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@+1 {{constant clauses expected}} - %5 = llvm.landingpad (catch %2 : !llvm.ptr) : !llvm.struct<(ptr, i32)> + %5 = llvm.landingpad (catch %2 : !llvm.ptr) : !llvm.struct<(ptr, i32)> llvm.return %0 : i32 } @@ -829,7 +782,7 @@ llvm.return %0 : i32 ^bb2: // pred: ^bb0 // expected-error@+1 {{landingpad instruction expects at least one clause or cleanup attribute}} - %2 = llvm.landingpad : !llvm.struct<(ptr, i32)> + %2 = llvm.landingpad : !llvm.struct<(ptr, i32)> llvm.return %0 : i32 } @@ -844,7 +797,7 @@ ^bb1: // pred: ^bb0 llvm.return %0 : i32 ^bb2: // pred: ^bb0 - %2 = llvm.landingpad cleanup : !llvm.struct<(ptr, i32)> + %2 = llvm.landingpad cleanup : !llvm.struct<(ptr, i32)> // expected-error@+1 {{'llvm.resume' op expects landingpad value as operand}} llvm.resume %0 : i32 } @@ -860,8 +813,8 @@ llvm.return %0 : i32 ^bb2: // pred: ^bb0 // expected-error@+1 {{llvm.landingpad needs to be in a function with a personality}} - %2 = llvm.landingpad cleanup : !llvm.struct<(ptr, i32)> - llvm.resume %2 : !llvm.struct<(ptr, i32)> + %2 = llvm.landingpad cleanup : !llvm.struct<(ptr, i32)> + llvm.resume %2 : !llvm.struct<(ptr, i32)> } // ----- @@ -1056,55 +1009,55 @@ // ----- -llvm.func @wmmaLoadOp_invalid_mem_space(%arg0: !llvm.ptr, %arg1: i32) { +llvm.func @wmmaLoadOp_invalid_mem_space(%arg0: !llvm.ptr<5>, %arg1: i32) { // expected-error@+1 {{'nvvm.wmma.load' op expected source pointer in memory space 0, 1, 3}} %0 = nvvm.wmma.load %arg0, %arg1 {eltype = #nvvm.mma_type, frag = #nvvm.mma_frag, k = 16 : i32, layout = #nvvm.mma_layout, m = 16 : i32, n = 16 : i32} - : (!llvm.ptr) -> !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, %arg1: i32) { +llvm.func @wmmaLoadOp_invalid_AOp(%arg0: !llvm.ptr<3>, %arg1: i32) { // expected-error@+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, frag = #nvvm.mma_frag, k = 16 : i32, layout = #nvvm.mma_layout, m = 16 : i32, n = 16 : i32} - : (!llvm.ptr) -> !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, %arg1: i32) { +llvm.func @wmmaLoadOp_invalid_BOp(%arg0: !llvm.ptr<3>, %arg1: i32) { // expected-error@+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, frag = #nvvm.mma_frag, k = 16 : i32, layout = #nvvm.mma_layout, m = 16 : i32, n = 16 : i32} - : (!llvm.ptr) -> !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, %arg1: i32) { +llvm.func @wmmaLoadOp_invalid_COp(%arg0: !llvm.ptr<3>, %arg1: i32) { // expected-error@+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, frag = #nvvm.mma_frag, k = 16 : i32, layout = #nvvm.mma_layout, m = 16 : i32, n = 16 : i32} - : (!llvm.ptr) -> !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, %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@+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, k = 16 : i32, layout = #nvvm.mma_layout, m = 16 : i32, n = 16 : i32} - : !llvm.ptr, 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 @wmmald_matrix(%arg0: !llvm.ptr) { +llvm.func @wmmald_matrix(%arg0: !llvm.ptr) { // expected-error@+1 {{'nvvm.ldmatrix' op expected source pointer in memory space 3}} - %l = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout} : (!llvm.ptr) -> i32 + %l = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout} : (!llvm.ptr) -> i32 llvm.return } // ----- -llvm.func @wmmald_matrix(%arg0: !llvm.ptr) { +llvm.func @wmmald_matrix(%arg0: !llvm.ptr<3>) { // expected-error@+1 {{'nvvm.ldmatrix' op expected num attribute to be 1, 2 or 4}} - %l = nvvm.ldmatrix %arg0 {num = 3 : i32, layout = #nvvm.mma_layout} : (!llvm.ptr) -> i32 + %l = nvvm.ldmatrix %arg0 {num = 3 : i32, layout = #nvvm.mma_layout} : (!llvm.ptr<3>) -> i32 llvm.return } // ----- -llvm.func @wmmald_matrix(%arg0: !llvm.ptr) { +llvm.func @wmmald_matrix(%arg0: !llvm.ptr<3>) { // expected-error@+1 {{'nvvm.ldmatrix' op expected destination type is i32}} - %l = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout} : (!llvm.ptr) -> !llvm.struct<(i32)> + %l = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout} : (!llvm.ptr<3>) -> !llvm.struct<(i32)> llvm.return } // ----- -llvm.func @wmmald_matrix(%arg0: !llvm.ptr) { +llvm.func @wmmald_matrix(%arg0: !llvm.ptr<3>) { // expected-error@+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} : (!llvm.ptr) -> !llvm.struct<(i32, i32)> + %l = nvvm.ldmatrix %arg0 {num = 4 : i32, layout = #nvvm.mma_layout} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32)> llvm.return } @@ -1278,33 +1231,33 @@ // ----- -func.func @cp_async(%arg0: !llvm.ptr, %arg1: !llvm.ptr) { +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, !llvm.ptr + nvvm.cp.async.shared.global %arg0, %arg1, 32 : !llvm.ptr<3>, !llvm.ptr<1> return } // ----- -func.func @cp_async(%arg0: !llvm.ptr, %arg1: !llvm.ptr) { +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, !llvm.ptr + 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>, %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>, i32, i32) -> !llvm.ptr + llvm.getelementptr %arg0[%arg1, %arg1] : (!llvm.ptr, i32, i32) -> !llvm.ptr, !llvm.struct<(i32)> return } // ----- -func.func @gep_out_of_bounds(%ptr: !llvm.ptr)>>, %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)>>, i64) -> !llvm.ptr + llvm.getelementptr %ptr[%idx, 1, 3] : (!llvm.ptr, i64) -> !llvm.ptr, !llvm.struct<(i32, struct<(i32, f32)>)> return } @@ -1321,8 +1274,8 @@ llvm.mlir.global internal @side_effecting_global() : !llvm.struct<(i8)> { %0 = llvm.mlir.constant(1 : i64) : i64 // expected-error@below {{ops with side effects not allowed in global initializers}} - %1 = llvm.alloca %0 x !llvm.struct<(i8)> : (i64) -> !llvm.ptr> - %2 = llvm.load %1 : !llvm.ptr> + %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 --- /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 + // CHECK: alignment = 8 + // CHECK: alloca_memory_space = 0 + // CHECK: bitsize = 64 + // CHECK: preferred = 8 + // CHECK: size = 8 + "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 + // CHECK: alignment = 8 + // CHECK: alloca_memory_space = 0 + // CHECK: bitsize = 64 + // CHECK: preferred = 8 + // CHECK: size = 8 + "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 + // CHECK: alignment = 8 + // CHECK: alloca_memory_space = 0 + // CHECK: bitsize = 64 + // CHECK: preferred = 8 + // CHECK: size = 8 + "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<5> + return + } +} + +// ----- + +module attributes { dlti.dl_spec = #dlti.dl_spec< + #dlti.dl_entry, dense<[32, 32, 64]> : vector<3xi32>>, + #dlti.dl_entry, dense<[64, 64, 64]> : vector<3xi32>>, + #dlti.dl_entry, 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 + // CHECK: alignment = 4 + // CHECK: alloca_memory_space = 5 + // CHECK: bitsize = 32 + // CHECK: preferred = 8 + // CHECK: size = 4 + "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 + // CHECK: alignment = 4 + // CHECK: alloca_memory_space = 5 + // CHECK: bitsize = 32 + // CHECK: preferred = 8 + // CHECK: size = 4 + "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 + // CHECK: alignment = 8 + // CHECK: alloca_memory_space = 5 + // CHECK: bitsize = 64 + // CHECK: preferred = 8 + // CHECK: size = 8 + "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<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@below {{unexpected layout attribute for pointer to 'i32'}} +module attributes { dlti.dl_spec = #dlti.dl_spec< + #dlti.dl_entry, dense<[64, 64, 64]> : vector<3xi32>> +>} { + func.func @pointer() { + return + } +} + +// ----- + +// expected-error@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, dense<[64.0, 64.0, 64.0]> : vector<3xf32>> +>} { + func.func @pointer() { + return + } +} + +// ----- + +// expected-error@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, 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 --- 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 - // CHECK: alignment = 8 - // CHECK: alloca_memory_space = 0 - // CHECK: bitsize = 64 - // CHECK: preferred = 8 - // CHECK: size = 8 - "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 - // CHECK: alignment = 8 - // CHECK: alloca_memory_space = 0 - // CHECK: bitsize = 64 - // CHECK: preferred = 8 - // CHECK: size = 8 - "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 - // CHECK: alignment = 8 - // CHECK: alloca_memory_space = 0 - // CHECK: bitsize = 64 - // CHECK: preferred = 8 - // CHECK: size = 8 - "test.data_layout_query"() : () -> !llvm.ptr + "test.data_layout_query"() : () -> !llvm.ptr<3> // CHECK: alignment = 8 // CHECK: alloca_memory_space = 0 // CHECK: bitsize = 64 @@ -52,8 +23,8 @@ // ----- module attributes { dlti.dl_spec = #dlti.dl_spec< - #dlti.dl_entry, dense<[32, 32, 64]> : vector<3xi32>>, - #dlti.dl_entry, dense<[64, 64, 64]> : vector<3xi32>>, + #dlti.dl_entry : vector<3xi32>>, + #dlti.dl_entry, dense<[64, 64, 64]> : vector<3xi32>>, #dlti.dl_entry, dense<[32, 64, 64]> : vector<3xi32>>, #dlti.dl_entry<"dlti.alloca_memory_space", 5 : ui32> >} { @@ -64,37 +35,19 @@ // CHECK: bitsize = 32 // CHECK: preferred = 8 // CHECK: size = 4 - "test.data_layout_query"() : () -> !llvm.ptr + "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 - // CHECK: alignment = 4 - // CHECK: alloca_memory_space = 5 - // CHECK: bitsize = 32 - // CHECK: preferred = 8 - // CHECK: size = 4 - "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> - // CHECK: alignment = 4 - // CHECK: alloca_memory_space = 5 - // CHECK: bitsize = 32 - // CHECK: preferred = 8 - // CHECK: size = 4 - "test.data_layout_query"() : () -> !llvm.ptr + "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 + "test.data_layout_query"() : () -> !llvm.ptr<5> // CHECK: alignment = 4 // CHECK: alloca_memory_space = 5 // CHECK: bitsize = 32 @@ -113,20 +66,9 @@ // ----- -// expected-error@below {{unexpected layout attribute for pointer to 'i32'}} -module attributes { dlti.dl_spec = #dlti.dl_spec< - #dlti.dl_entry, dense<[64, 64, 64]> : vector<3xi32>> ->} { - func.func @pointer() { - return - } -} - -// ----- - -// expected-error@below {{expected layout attribute for '!llvm.ptr' to be a dense integer elements attribute with 3 or 4 elements}} +// expected-error@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, dense<[64.0, 64.0, 64.0]> : vector<3xf32>> + #dlti.dl_entry : vector<3xf32>> >} { func.func @pointer() { return @@ -137,7 +79,7 @@ // expected-error@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, dense<[64, 64, 32]> : vector<3xi32>> + #dlti.dl_entry : 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 --- /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, %arg1 : i32) -> !llvm.struct<(i32, i32, i32, i32)> { + // CHECK: nvvm.wmma.load {{.*}} {eltype = #nvvm.mma_type, frag = #nvvm.mma_frag, k = 8 : i32, layout = #nvvm.mma_layout, m = 16 : i32, n = 16 : i32} + %0 = nvvm.wmma.load %arg0, %arg1 + {eltype = #nvvm.mma_type, frag = #nvvm.mma_frag, k = 8 : i32, layout = #nvvm.mma_layout, m = 16 : i32, n = 16 : i32} + : (!llvm.ptr) -> !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, %arg1: !llvm.ptr) { +// CHECK: nvvm.cp.async.shared.global %{{.*}}, %{{.*}}, 16 + nvvm.cp.async.shared.global %arg0, %arg1, 16 : !llvm.ptr, !llvm.ptr +// CHECK: nvvm.cp.async.shared.global %{{.*}}, %{{.*}}, 16 {bypass_l1} + nvvm.cp.async.shared.global %arg0, %arg1, 16 {bypass_l1} : !llvm.ptr, !llvm.ptr +// 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) { + // CHECK: nvvm.ldmatrix %{{.*}} {layout = #nvvm.mma_layout, num = 1 : i32} : (!llvm.ptr) -> i32 + %l1 = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout} : (!llvm.ptr) -> i32 + // CHECK: nvvm.ldmatrix %{{.*}} {layout = #nvvm.mma_layout, num = 2 : i32} : (!llvm.ptr) -> !llvm.struct<(i32, i32)> + %l2 = nvvm.ldmatrix %arg0 {num = 2 : i32, layout = #nvvm.mma_layout} : (!llvm.ptr) -> !llvm.struct<(i32, i32)> + // CHECK: nvvm.ldmatrix %{{.*}} {layout = #nvvm.mma_layout, num = 4 : i32} : (!llvm.ptr) -> !llvm.struct<(i32, i32, i32, i32)> + %l4 = nvvm.ldmatrix %arg0 {num = 4 : i32, layout = #nvvm.mma_layout} : (!llvm.ptr) -> !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 --- a/mlir/test/Dialect/LLVMIR/nvvm.mlir +++ b/mlir/test/Dialect/LLVMIR/nvvm.mlir @@ -266,11 +266,11 @@ } // CHECK-LABEL: @nvvm_wmma_load_tf32 -func.func @nvvm_wmma_load_tf32(%arg0: !llvm.ptr, %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, frag = #nvvm.mma_frag, k = 8 : i32, layout = #nvvm.mma_layout, m = 16 : i32, n = 16 : i32} %0 = nvvm.wmma.load %arg0, %arg1 {eltype = #nvvm.mma_type, frag = #nvvm.mma_frag, k = 8 : i32, layout = #nvvm.mma_layout, m = 16 : i32, n = 16 : i32} - : (!llvm.ptr) -> !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 @@ } // CHECK-LABEL: @cp_async -llvm.func @cp_async(%arg0: !llvm.ptr, %arg1: !llvm.ptr) { +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, !llvm.ptr + 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, !llvm.ptr + 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 @@ } // CHECK-LABEL: llvm.func @ld_matrix -llvm.func @ld_matrix(%arg0: !llvm.ptr) { - // CHECK: nvvm.ldmatrix %{{.*}} {layout = #nvvm.mma_layout, num = 1 : i32} : (!llvm.ptr) -> i32 - %l1 = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout} : (!llvm.ptr) -> i32 - // CHECK: nvvm.ldmatrix %{{.*}} {layout = #nvvm.mma_layout, num = 2 : i32} : (!llvm.ptr) -> !llvm.struct<(i32, i32)> - %l2 = nvvm.ldmatrix %arg0 {num = 2 : i32, layout = #nvvm.mma_layout} : (!llvm.ptr) -> !llvm.struct<(i32, i32)> - // CHECK: nvvm.ldmatrix %{{.*}} {layout = #nvvm.mma_layout, num = 4 : i32} : (!llvm.ptr) -> !llvm.struct<(i32, i32, i32, i32)> - %l4 = nvvm.ldmatrix %arg0 {num = 4 : i32, layout = #nvvm.mma_layout} : (!llvm.ptr) -> !llvm.struct<(i32, i32, i32, i32)> +llvm.func @ld_matrix(%arg0: !llvm.ptr<3>) { + // CHECK: nvvm.ldmatrix %{{.*}} {layout = #nvvm.mma_layout, num = 1 : i32} : (!llvm.ptr<3>) -> i32 + %l1 = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout} : (!llvm.ptr<3>) -> i32 + // CHECK: nvvm.ldmatrix %{{.*}} {layout = #nvvm.mma_layout, num = 2 : i32} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32)> + %l2 = nvvm.ldmatrix %arg0 {num = 2 : i32, layout = #nvvm.mma_layout} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32)> + // CHECK: nvvm.ldmatrix %{{.*}} {layout = #nvvm.mma_layout, num = 4 : i32} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32, i32, i32)> + %l4 = nvvm.ldmatrix %arg0 {num = 4 : i32, layout = #nvvm.mma_layout} : (!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 @@ // 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 --- /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@below {{"llvm.sret" attribute attached to LLVM pointer argument of different type}} +llvm.func @invalid_sret_attr_type(%0 : !llvm.ptr {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 --- a/mlir/test/Dialect/LLVMIR/parameter-attrs-invalid.mlir +++ b/mlir/test/Dialect/LLVMIR/parameter-attrs-invalid.mlir @@ -47,11 +47,6 @@ // ----- -// expected-error@below {{"llvm.sret" attribute attached to LLVM pointer argument of different type}} -llvm.func @invalid_sret_attr_type(%0 : !llvm.ptr {llvm.sret = !llvm.struct<(i32)>}) - -// ----- - // expected-error@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 --- /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 +} + +// ----- + +func.func @repeated_struct_name() { + "some.op"() : () -> !llvm.struct<"a", (ptr>)> + // expected-error @+1 {{identified type already used with a different body}} + "some.op"() : () -> !llvm.struct<"a", (i32)> +} + +// ----- + +func.func @dynamic_vector() { + // expected-error @+1 {{expected '? x x ' or ' x '}} + "some.op"() : () -> !llvm.vec> +} + +// ----- + +func.func @dynamic_scalable_vector() { + // expected-error @+1 {{expected '? x x ' or ' x '}} + "some.op"() : () -> !llvm.vec> +} + +// ----- + +func.func @unscalable_vector() { + // expected-error @+1 {{expected '? x x ' or ' x '}} + "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> +} diff --git a/mlir/test/Dialect/LLVMIR/types-invalid.mlir b/mlir/test/Dialect/LLVMIR/types-invalid.mlir --- a/mlir/test/Dialect/LLVMIR/types-invalid.mlir +++ b/mlir/test/Dialect/LLVMIR/types-invalid.mlir @@ -21,15 +21,8 @@ // ----- -func.func @void_pointer() { - // expected-error @+1 {{invalid pointer element type}} - "some.op"() : () -> !llvm.ptr -} - -// ----- - func.func @repeated_struct_name() { - "some.op"() : () -> !llvm.struct<"a", (ptr>)> + "some.op"() : () -> !llvm.struct<"a", (ptr)> // expected-error @+1 {{identified type already used with a different body}} "some.op"() : () -> !llvm.struct<"a", (i32)> } @@ -113,28 +106,28 @@ func.func @dynamic_vector() { // expected-error @+1 {{expected '? x x ' or ' x '}} - "some.op"() : () -> !llvm.vec> + "some.op"() : () -> !llvm.vec } // ----- func.func @dynamic_scalable_vector() { // expected-error @+1 {{expected '? x x ' or ' x '}} - "some.op"() : () -> !llvm.vec> + "some.op"() : () -> !llvm.vec } // ----- func.func @unscalable_vector() { // expected-error @+1 {{expected '? x x ' or ' x '}} - "some.op"() : () -> !llvm.vec<4x4 x ptr> + "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> + "some.op"() : () -> !llvm.vec<0 x ptr> } // ----- diff --git a/mlir/test/Dialect/LLVMIR/types.mlir b/mlir/test/Dialect/LLVMIR/types-typed-pointers.mlir copy from mlir/test/Dialect/LLVMIR/types.mlir copy to mlir/test/Dialect/LLVMIR/types-typed-pointers.mlir --- a/mlir/test/Dialect/LLVMIR/types.mlir +++ b/mlir/test/Dialect/LLVMIR/types-typed-pointers.mlir @@ -1,60 +1,5 @@ // RUN: mlir-opt -allow-unregistered-dialect %s -split-input-file | mlir-opt -allow-unregistered-dialect | FileCheck %s -// CHECK-LABEL: @primitive -func.func @primitive() { - // CHECK: !llvm.void - "some.op"() : () -> !llvm.void - // CHECK: !llvm.ppc_fp128 - "some.op"() : () -> !llvm.ppc_fp128 - // CHECK: !llvm.x86_mmx - "some.op"() : () -> !llvm.x86_mmx - // CHECK: !llvm.token - "some.op"() : () -> !llvm.token - // CHECK: !llvm.label - "some.op"() : () -> !llvm.label - // CHECK: !llvm.metadata - "some.op"() : () -> !llvm.metadata - return -} - -// CHECK-LABEL: @func -func.func @func() { - // CHECK: !llvm.func - "some.op"() : () -> !llvm.func - // CHECK: !llvm.func - "some.op"() : () -> !llvm.func - // CHECK: !llvm.func - "some.op"() : () -> !llvm.func - // CHECK: !llvm.func - "some.op"() : () -> !llvm.func - // CHECK: !llvm.func - "some.op"() : () -> !llvm.func - // CHECK: !llvm.func - "some.op"() : () -> !llvm.func - // CHECK: !llvm.func - "some.op"() : () -> !llvm.func - return -} - -// CHECK-LABEL: @integer -func.func @integer() { - // CHECK: i1 - "some.op"() : () -> i1 - // CHECK: i8 - "some.op"() : () -> i8 - // CHECK: i16 - "some.op"() : () -> i16 - // CHECK: i32 - "some.op"() : () -> i32 - // CHECK: i64 - "some.op"() : () -> i64 - // CHECK: i57 - "some.op"() : () -> i57 - // CHECK: i129 - "some.op"() : () -> i129 - return -} - // CHECK-LABEL: @ptr func.func @ptr() { // CHECK: !llvm.ptr @@ -108,39 +53,6 @@ return } -// CHECK-LABEL: @literal_struct -func.func @literal_struct() { - // CHECK: !llvm.struct<()> - "some.op"() : () -> !llvm.struct<()> - // CHECK: !llvm.struct<(i32)> - "some.op"() : () -> !llvm.struct<(i32)> - // CHECK: !llvm.struct<(f32, i32)> - "some.op"() : () -> !llvm.struct<(f32, i32)> - // CHECK: !llvm.struct<(struct<(i32)>)> - "some.op"() : () -> !llvm.struct<(struct<(i32)>)> - // CHECK: !llvm.struct<(i32, struct<(i32)>, f32)> - "some.op"() : () -> !llvm.struct<(i32, struct<(i32)>, f32)> - - // CHECK: !llvm.struct - "some.op"() : () -> !llvm.struct - // CHECK: !llvm.struct - "some.op"() : () -> !llvm.struct - // CHECK: !llvm.struct - "some.op"() : () -> !llvm.struct - // CHECK: !llvm.struct - "some.op"() : () -> !llvm.struct - // CHECK: !llvm.struct)> - "some.op"() : () -> !llvm.struct)> - // CHECK: !llvm.struct, f32)> - "some.op"() : () -> !llvm.struct, f32)> - - // CHECK: !llvm.struct<(struct)> - "some.op"() : () -> !llvm.struct<(struct)> - // CHECK: !llvm.struct)> - "some.op"() : () -> !llvm.struct)> - return -} - // CHECK-LABEL: @identified_struct func.func @identified_struct() { // CHECK: !llvm.struct<"empty", ()> @@ -174,12 +86,6 @@ return } -func.func @verbose() { - // CHECK: !llvm.struct<(i64, struct<(f32)>)> - "some.op"() : () -> !llvm.struct<(i64, !llvm.struct<(f32)>)> - return -} - // CHECK-LABEL: @ptr_elem_interface // CHECK-COUNT-3: !llvm.ptr // CHECK: llvm.mlir.undef : !llvm.ptr diff --git a/mlir/test/Dialect/LLVMIR/types.mlir b/mlir/test/Dialect/LLVMIR/types.mlir --- a/mlir/test/Dialect/LLVMIR/types.mlir +++ b/mlir/test/Dialect/LLVMIR/types.mlir @@ -57,26 +57,14 @@ // CHECK-LABEL: @ptr func.func @ptr() { - // CHECK: !llvm.ptr - "some.op"() : () -> !llvm.ptr - // CHECK: !llvm.ptr - "some.op"() : () -> !llvm.ptr - // CHECK: !llvm.ptr> - "some.op"() : () -> !llvm.ptr> - // CHECK: !llvm.ptr>>>> - "some.op"() : () -> !llvm.ptr>>>> - // CHECK: !llvm.ptr - "some.op"() : () -> !llvm.ptr - // CHECK: !llvm.ptr - "some.op"() : () -> !llvm.ptr - // CHECK: !llvm.ptr - "some.op"() : () -> !llvm.ptr - // CHECK: !llvm.ptr, 9> - "some.op"() : () -> !llvm.ptr, 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, 9> + "some.op"() : () -> !llvm.ptr, 9> return } @@ -90,8 +78,8 @@ "some.op"() : () -> !llvm.vec // CHECK: !llvm.vec "some.op"() : () -> !llvm.vec - // CHECK: !llvm.vec<4 x ptr> - "some.op"() : () -> !llvm.vec<4 x ptr> + // CHECK: !llvm.vec<4 x ptr> + "some.op"() : () -> !llvm.vec<4 x ptr> return } @@ -101,8 +89,8 @@ "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> - "some.op"() : () -> !llvm.array<10 x ptr> + // 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 @@ "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>)> - "some.op"() : () -> !llvm.struct<"long", (i32, struct<(i32, i1)>, f32, ptr>)> - // CHECK: !llvm.struct<"self-recursive", (ptr>)> - "some.op"() : () -> !llvm.struct<"self-recursive", (ptr>)> + // 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, 3>)>>)> - "some.op"() : () -> !llvm.struct<"mutually-a", (ptr, 3>)>>)> - // CHECK: !llvm.struct<"mutually-b", (ptr>)>, 3>)> - "some.op"() : () -> !llvm.struct<"mutually-b", (ptr>)>, 3>)> - // CHECK: !llvm.struct<"referring-another", (ptr>)> - "some.op"() : () -> !llvm.struct<"referring-another", (ptr>)> - + // 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> - "some.op"() : () -> !llvm.ptr> return } @@ -180,16 +160,6 @@ return } -// CHECK-LABEL: @ptr_elem_interface -// CHECK-COUNT-3: !llvm.ptr -// CHECK: llvm.mlir.undef : !llvm.ptr -func.func @ptr_elem_interface(%arg0: !llvm.ptr) { - %0 = llvm.load %arg0 : !llvm.ptr - llvm.store %0, %arg0 : !llvm.ptr - llvm.mlir.undef : !llvm.ptr - return -} - // ----- // Check that type aliases can be used inside LLVM dialect types. Note that @@ -200,13 +170,9 @@ !baz = i64 !qux = !llvm.struct<(!baz)> -!rec = !llvm.struct<"a", (ptr>)> - // CHECK: aliases llvm.func @aliases() { // CHECK: !llvm.struct<(i32, f32, struct<(i64)>)> "some.op"() : () -> !llvm.struct<(i32, f32, !qux)> - // CHECK: !llvm.struct<"a", (ptr>)> - "some.op"() : () -> !rec llvm.return }