[Mlir-commits] [mlir] [MLIR][LLVM] Remove last typed pointer remnants from tests (PR #71232)
Christian Ulmann
llvmlistbot at llvm.org
Fri Nov 3 14:18:35 PDT 2023
https://github.com/Dinistro created https://github.com/llvm/llvm-project/pull/71232
This commit removes all LLVM dialect typed pointers from the lit tests. Typed pointers have been deprecated for a while now and it's planned to soon remove them from the LLVM dialect.
Related PSA: https://discourse.llvm.org/t/psa-removal-of-typed-pointers-from-the-llvm-dialect/74502
>From 159cecee1b76ec1119d0969a2e8a427d94ccc00c Mon Sep 17 00:00:00 2001
From: Christian Ulmann <christianulmann at gmail.com>
Date: Fri, 3 Nov 2023 22:15:27 +0100
Subject: [PATCH] [MLIR][LLVM] Remove last typed pointer remnants from tests
This commit removes all LLVM dialect typed pointers from the lit tests.
Typed pointers have been deprecated for a while now and it's planned to
soon remove them from the LLVM dialect.
Related PSA: https://discourse.llvm.org/t/psa-removal-of-typed-pointers-from-the-llvm-dialect/74502
---
.../convert-types-typed-pointers.mlir | 40 ---
...ower-launch-func-to-gpu-runtime-calls.mlir | 8 +-
.../OpenMPToLLVM/convert-to-llvmir.mlir | 180 ++++++-----
mlir/test/Dialect/GPU/invalid.mlir | 20 +-
.../Dialect/GPU/module-to-binary-nvvm.mlir | 8 +-
.../Dialect/GPU/module-to-binary-rocdl.mlir | 8 +-
mlir/test/Dialect/GPU/outlining.mlir | 4 +-
.../LLVMIR/canonicalize-typed-pointers.mlir | 86 -----
.../dynamic-gep-index-typed-pointers.mlir | 12 -
.../Dialect/LLVMIR/global-typed-pointers.mlir | 46 ---
.../LLVMIR/invalid-typed-pointers.mlir | 299 ------------------
.../Dialect/LLVMIR/layout-typed-pointers.mlir | 145 ---------
.../Dialect/LLVMIR/nvvm-typed-pointers.mlir | 55 ----
...arameter-attrs-invalid-typed-pointers.mlir | 6 -
.../LLVMIR/parameter-attrs-invalid.mlir | 15 -
.../LLVMIR/roundtrip-typed-pointers.mlir | 73 -----
mlir/test/Dialect/LLVMIR/sroa.mlir | 14 -
.../LLVMIR/types-invalid-typed-pointers.mlir | 42 ---
.../Dialect/LLVMIR/types-typed-pointers.mlir | 118 -------
mlir/test/Dialect/LLVMIR/types.mlir | 4 +-
mlir/test/Dialect/OpenACC/invalid.mlir | 72 ++---
mlir/test/Dialect/OpenACC/ops.mlir | 280 ++++++++--------
mlir/test/Dialect/OpenMP/canonicalize.mlir | 4 +-
mlir/test/Dialect/OpenMP/invalid.mlir | 84 ++---
mlir/test/Dialect/OpenMP/ops.mlir | 172 +++++-----
mlir/test/Dialect/SparseTensor/invalid.mlir | 8 +-
..._shot_bufferize_tensor_copy_insertion.mlir | 2 +-
.../SparseTensor/rewriting_for_codegen.mlir | 30 +-
mlir/test/Dialect/SparseTensor/roundtrip.mlir | 16 +-
mlir/test/mlir-cpu-runner/simple.mlir | 28 +-
mlir/test/mlir-cpu-runner/x86-varargs.mlir | 36 +--
.../dialects/gpu/module-to-binary-nvvm.py | 8 +-
.../dialects/gpu/module-to-binary-rocdl.py | 8 +-
33 files changed, 485 insertions(+), 1446 deletions(-)
delete mode 100644 mlir/test/Conversion/FuncToLLVM/convert-types-typed-pointers.mlir
delete mode 100644 mlir/test/Dialect/LLVMIR/canonicalize-typed-pointers.mlir
delete mode 100644 mlir/test/Dialect/LLVMIR/dynamic-gep-index-typed-pointers.mlir
delete mode 100644 mlir/test/Dialect/LLVMIR/global-typed-pointers.mlir
delete mode 100644 mlir/test/Dialect/LLVMIR/invalid-typed-pointers.mlir
delete mode 100644 mlir/test/Dialect/LLVMIR/layout-typed-pointers.mlir
delete mode 100644 mlir/test/Dialect/LLVMIR/nvvm-typed-pointers.mlir
delete mode 100644 mlir/test/Dialect/LLVMIR/parameter-attrs-invalid-typed-pointers.mlir
delete mode 100644 mlir/test/Dialect/LLVMIR/roundtrip-typed-pointers.mlir
delete mode 100644 mlir/test/Dialect/LLVMIR/types-invalid-typed-pointers.mlir
delete mode 100644 mlir/test/Dialect/LLVMIR/types-typed-pointers.mlir
diff --git a/mlir/test/Conversion/FuncToLLVM/convert-types-typed-pointers.mlir b/mlir/test/Conversion/FuncToLLVM/convert-types-typed-pointers.mlir
deleted file mode 100644
index c55bb8bb71c676b..000000000000000
--- a/mlir/test/Conversion/FuncToLLVM/convert-types-typed-pointers.mlir
+++ /dev/null
@@ -1,40 +0,0 @@
-// RUN: mlir-opt -test-convert-call-op %s | FileCheck %s
-
-// CHECK-LABEL: @ptr
-// CHECK: !llvm.ptr<i42>
-func.func private @ptr() -> !llvm.ptr<!test.smpla>
-
-// CHECK-LABEL: @opaque_ptr
-// CHECK: !llvm.ptr
-// CHECK-NOT: <
-func.func private @opaque_ptr() -> !llvm.ptr
-
-// CHECK-LABEL: @ptr_ptr()
-// CHECK: !llvm.ptr<ptr<i42>>
-func.func private @ptr_ptr() -> !llvm.ptr<!llvm.ptr<!test.smpla>>
-
-// CHECK-LABEL: @struct_ptr()
-// CHECK: !llvm.struct<(ptr<i42>)>
-func.func private @struct_ptr() -> !llvm.struct<(ptr<!test.smpla>)>
-
-// CHECK-LABEL: @named_struct_ptr()
-// CHECK: !llvm.struct<"_Converted_named", (ptr<i42>)>
-func.func private @named_struct_ptr() -> !llvm.struct<"named", (ptr<!test.smpla>)>
-
-// CHECK-LABEL: @named_no_convert
-// CHECK: !llvm.struct<"no_convert", (ptr<struct<"no_convert">>)>
-func.func private @named_no_convert() -> !llvm.struct<"no_convert", (ptr<struct<"no_convert">>)>
-
-// CHECK-LABEL: @array_ptr()
-// CHECK: !llvm.array<10 x ptr<i42>>
-func.func private @array_ptr() -> !llvm.array<10 x ptr<!test.smpla>>
-
-// CHECK-LABEL: @func()
-// CHECK: !llvm.ptr<func<i42 (i42)>>
-func.func private @func() -> !llvm.ptr<!llvm.func<!test.smpla (!test.smpla)>>
-
-// TODO: support conversion of recursive types in the conversion infra.
-// CHECK-LABEL: @named_recursive()
-// CHECK: !llvm.struct<"_Converted_recursive", (ptr<i42>, ptr<struct<"_Converted_recursive">>)>
-func.func private @named_recursive() -> !llvm.struct<"recursive", (ptr<!test.smpla>, ptr<struct<"recursive">>)>
-
diff --git a/mlir/test/Conversion/GPUCommon/lower-launch-func-to-gpu-runtime-calls.mlir b/mlir/test/Conversion/GPUCommon/lower-launch-func-to-gpu-runtime-calls.mlir
index 1b9afcdf50a17f0..f5462b579b5eb0c 100644
--- a/mlir/test/Conversion/GPUCommon/lower-launch-func-to-gpu-runtime-calls.mlir
+++ b/mlir/test/Conversion/GPUCommon/lower-launch-func-to-gpu-runtime-calls.mlir
@@ -10,8 +10,8 @@ module attributes {gpu.container_module} {
gpu.module @kernel_module attributes {
nvvm.cubin = "CUBIN", rocdl.hsaco = "HSACO"
} {
- llvm.func @kernel(%arg0: i32, %arg1: !llvm.ptr<f32>,
- %arg2: !llvm.ptr<f32>, %arg3: i64, %arg4: i64,
+ llvm.func @kernel(%arg0: i32, %arg1: !llvm.ptr,
+ %arg2: !llvm.ptr, %arg3: i64, %arg4: i64,
%arg5: i64) attributes {gpu.kernel} {
llvm.return
}
@@ -69,8 +69,8 @@ module attributes {gpu.container_module} {
// CHECK: gpu.module
// ROCDL: gpu.module
gpu.module @kernel_module [#nvvm.target] {
- llvm.func @kernel(%arg0: i32, %arg1: !llvm.ptr<f32>,
- %arg2: !llvm.ptr<f32>, %arg3: i64, %arg4: i64,
+ llvm.func @kernel(%arg0: i32, %arg1: !llvm.ptr,
+ %arg2: !llvm.ptr, %arg3: i64, %arg4: i64,
%arg5: i64) attributes {gpu.kernel} {
llvm.return
}
diff --git a/mlir/test/Conversion/OpenMPToLLVM/convert-to-llvmir.mlir b/mlir/test/Conversion/OpenMPToLLVM/convert-to-llvmir.mlir
index 9c777703bb73970..bbf50617edf9448 100644
--- a/mlir/test/Conversion/OpenMPToLLVM/convert-to-llvmir.mlir
+++ b/mlir/test/Conversion/OpenMPToLLVM/convert-to-llvmir.mlir
@@ -88,30 +88,30 @@ func.func @wsloop(%arg0: index, %arg1: index, %arg2: index, %arg3: index, %arg4:
// -----
// CHECK-LABEL: @atomic_write
-// CHECK: (%[[ARG0:.*]]: !llvm.ptr<i32>)
+// CHECK: (%[[ARG0:.*]]: !llvm.ptr)
// CHECK: %[[VAL0:.*]] = llvm.mlir.constant(1 : i32) : i32
-// CHECK: omp.atomic.write %[[ARG0]] = %[[VAL0]] memory_order(relaxed) : !llvm.ptr<i32>, i32
-func.func @atomic_write(%a: !llvm.ptr<i32>) -> () {
+// CHECK: omp.atomic.write %[[ARG0]] = %[[VAL0]] memory_order(relaxed) : !llvm.ptr, i32
+func.func @atomic_write(%a: !llvm.ptr) -> () {
%1 = arith.constant 1 : i32
- omp.atomic.write %a = %1 hint(none) memory_order(relaxed) : !llvm.ptr<i32>, i32
+ omp.atomic.write %a = %1 hint(none) memory_order(relaxed) : !llvm.ptr, i32
return
}
// -----
// CHECK-LABEL: @atomic_read
-// CHECK: (%[[ARG0:.*]]: !llvm.ptr<i32>, %[[ARG1:.*]]: !llvm.ptr<i32>)
-// CHECK: omp.atomic.read %[[ARG1]] = %[[ARG0]] memory_order(acquire) hint(contended) : !llvm.ptr<i32>
-func.func @atomic_read(%a: !llvm.ptr<i32>, %b: !llvm.ptr<i32>) -> () {
- omp.atomic.read %b = %a memory_order(acquire) hint(contended) : !llvm.ptr<i32>, i32
+// CHECK: (%[[ARG0:.*]]: !llvm.ptr, %[[ARG1:.*]]: !llvm.ptr)
+// CHECK: omp.atomic.read %[[ARG1]] = %[[ARG0]] memory_order(acquire) hint(contended) : !llvm.ptr
+func.func @atomic_read(%a: !llvm.ptr, %b: !llvm.ptr) -> () {
+ omp.atomic.read %b = %a memory_order(acquire) hint(contended) : !llvm.ptr, i32
return
}
// -----
func.func @atomic_update() {
- %0 = llvm.mlir.addressof @_QFsEc : !llvm.ptr<i32>
- omp.atomic.update %0 : !llvm.ptr<i32> {
+ %0 = llvm.mlir.addressof @_QFsEc : !llvm.ptr
+ omp.atomic.update %0 : !llvm.ptr {
^bb0(%arg0: i32):
%1 = arith.constant 1 : i32
%2 = arith.addi %arg0, %1 : i32
@@ -125,8 +125,8 @@ llvm.mlir.global internal @_QFsEc() : i32 {
}
// CHECK-LABEL: @atomic_update
-// CHECK: %[[GLOBAL_VAR:.*]] = llvm.mlir.addressof @_QFsEc : !llvm.ptr<i32>
-// CHECK: omp.atomic.update %[[GLOBAL_VAR]] : !llvm.ptr<i32> {
+// CHECK: %[[GLOBAL_VAR:.*]] = llvm.mlir.addressof @_QFsEc : !llvm.ptr
+// CHECK: omp.atomic.update %[[GLOBAL_VAR]] : !llvm.ptr {
// CHECK: ^bb0(%[[IN_VAL:.*]]: i32):
// CHECK: %[[CONST_1:.*]] = llvm.mlir.constant(1 : i32) : i32
// CHECK: %[[OUT_VAL:.*]] = llvm.add %[[IN_VAL]], %[[CONST_1]] : i32
@@ -136,10 +136,10 @@ llvm.mlir.global internal @_QFsEc() : i32 {
// -----
// CHECK-LABEL: @threadprivate
-// CHECK: (%[[ARG0:.*]]: !llvm.ptr<i32>)
-// CHECK: %[[VAL0:.*]] = omp.threadprivate %[[ARG0]] : !llvm.ptr<i32> -> !llvm.ptr<i32>
-func.func @threadprivate(%a: !llvm.ptr<i32>) -> () {
- %1 = omp.threadprivate %a : !llvm.ptr<i32> -> !llvm.ptr<i32>
+// CHECK: (%[[ARG0:.*]]: !llvm.ptr)
+// CHECK: %[[VAL0:.*]] = omp.threadprivate %[[ARG0]] : !llvm.ptr -> !llvm.ptr
+func.func @threadprivate(%a: !llvm.ptr) -> () {
+ %1 = omp.threadprivate %a : !llvm.ptr -> !llvm.ptr
return
}
@@ -175,15 +175,15 @@ func.func @simdloop_block_arg(%val : i32, %ub : i32, %i : index) {
// -----
// CHECK-LABEL: @task_depend
-// CHECK: (%[[ARG0:.*]]: !llvm.ptr<i32>) {
-// CHECK: omp.task depend(taskdependin -> %[[ARG0]] : !llvm.ptr<i32>) {
+// CHECK: (%[[ARG0:.*]]: !llvm.ptr) {
+// CHECK: omp.task depend(taskdependin -> %[[ARG0]] : !llvm.ptr) {
// CHECK: omp.terminator
// CHECK: }
// CHECK: llvm.return
// CHECK: }
-func.func @task_depend(%arg0: !llvm.ptr<i32>) {
- omp.task depend(taskdependin -> %arg0 : !llvm.ptr<i32>) {
+func.func @task_depend(%arg0: !llvm.ptr) {
+ omp.task depend(taskdependin -> %arg0 : !llvm.ptr) {
omp.terminator
}
return
@@ -192,47 +192,47 @@ func.func @task_depend(%arg0: !llvm.ptr<i32>) {
// -----
// CHECK-LABEL: @_QPomp_target_data
-// CHECK: (%[[ARG0:.*]]: !llvm.ptr<i32>, %[[ARG1:.*]]: !llvm.ptr<i32>, %[[ARG2:.*]]: !llvm.ptr<i32>, %[[ARG3:.*]]: !llvm.ptr<i32>)
-// CHECK: %[[MAP0:.*]] = omp.map_info var_ptr(%[[ARG0]] : !llvm.ptr<i32>, i32) map_clauses(to) capture(ByRef) -> !llvm.ptr<i32> {name = ""}
-// CHECK: %[[MAP1:.*]] = omp.map_info var_ptr(%[[ARG1]] : !llvm.ptr<i32>, i32) map_clauses(to) capture(ByRef) -> !llvm.ptr<i32> {name = ""}
-// CHECK: %[[MAP2:.*]] = omp.map_info var_ptr(%[[ARG2]] : !llvm.ptr<i32>, i32) map_clauses(always, exit_release_or_enter_alloc) capture(ByRef) -> !llvm.ptr<i32> {name = ""}
-// CHECK: omp.target_enter_data map_entries(%[[MAP0]], %[[MAP1]], %[[MAP2]] : !llvm.ptr<i32>, !llvm.ptr<i32>, !llvm.ptr<i32>)
-// CHECK: %[[MAP3:.*]] = omp.map_info var_ptr(%[[ARG0]] : !llvm.ptr<i32>, i32) map_clauses(from) capture(ByRef) -> !llvm.ptr<i32> {name = ""}
-// CHECK: %[[MAP4:.*]] = omp.map_info var_ptr(%[[ARG1]] : !llvm.ptr<i32>, i32) map_clauses(from) capture(ByRef) -> !llvm.ptr<i32> {name = ""}
-// CHECK: %[[MAP5:.*]] = omp.map_info var_ptr(%[[ARG2]] : !llvm.ptr<i32>, i32) map_clauses(exit_release_or_enter_alloc) capture(ByRef) -> !llvm.ptr<i32> {name = ""}
-// CHECK: %[[MAP6:.*]] = omp.map_info var_ptr(%[[ARG3]] : !llvm.ptr<i32>, i32) map_clauses(always, delete) capture(ByRef) -> !llvm.ptr<i32> {name = ""}
-// CHECK: omp.target_exit_data map_entries(%[[MAP3]], %[[MAP4]], %[[MAP5]], %[[MAP6]] : !llvm.ptr<i32>, !llvm.ptr<i32>, !llvm.ptr<i32>, !llvm.ptr<i32>)
-
-llvm.func @_QPomp_target_data(%a : !llvm.ptr<i32>, %b : !llvm.ptr<i32>, %c : !llvm.ptr<i32>, %d : !llvm.ptr<i32>) {
- %0 = omp.map_info var_ptr(%a : !llvm.ptr<i32>, i32) map_clauses(to) capture(ByRef) -> !llvm.ptr<i32> {name = ""}
- %1 = omp.map_info var_ptr(%b : !llvm.ptr<i32>, i32) map_clauses(to) capture(ByRef) -> !llvm.ptr<i32> {name = ""}
- %2 = omp.map_info var_ptr(%c : !llvm.ptr<i32>, i32) map_clauses(always, exit_release_or_enter_alloc) capture(ByRef) -> !llvm.ptr<i32> {name = ""}
- omp.target_enter_data map_entries(%0, %1, %2 : !llvm.ptr<i32>, !llvm.ptr<i32>, !llvm.ptr<i32>) {}
- %3 = omp.map_info var_ptr(%a : !llvm.ptr<i32>, i32) map_clauses(from) capture(ByRef) -> !llvm.ptr<i32> {name = ""}
- %4 = omp.map_info var_ptr(%b : !llvm.ptr<i32>, i32) map_clauses(from) capture(ByRef) -> !llvm.ptr<i32> {name = ""}
- %5 = omp.map_info var_ptr(%c : !llvm.ptr<i32>, i32) map_clauses(exit_release_or_enter_alloc) capture(ByRef) -> !llvm.ptr<i32> {name = ""}
- %6 = omp.map_info var_ptr(%d : !llvm.ptr<i32>, i32) map_clauses(always, delete) capture(ByRef) -> !llvm.ptr<i32> {name = ""}
- omp.target_exit_data map_entries(%3, %4, %5, %6 : !llvm.ptr<i32>, !llvm.ptr<i32>, !llvm.ptr<i32>, !llvm.ptr<i32>) {}
+// CHECK: (%[[ARG0:.*]]: !llvm.ptr, %[[ARG1:.*]]: !llvm.ptr, %[[ARG2:.*]]: !llvm.ptr, %[[ARG3:.*]]: !llvm.ptr)
+// CHECK: %[[MAP0:.*]] = omp.map_info var_ptr(%[[ARG0]] : !llvm.ptr, i32) map_clauses(to) capture(ByRef) -> !llvm.ptr {name = ""}
+// CHECK: %[[MAP1:.*]] = omp.map_info var_ptr(%[[ARG1]] : !llvm.ptr, i32) map_clauses(to) capture(ByRef) -> !llvm.ptr {name = ""}
+// CHECK: %[[MAP2:.*]] = omp.map_info var_ptr(%[[ARG2]] : !llvm.ptr, i32) map_clauses(always, exit_release_or_enter_alloc) capture(ByRef) -> !llvm.ptr {name = ""}
+// CHECK: omp.target_enter_data map_entries(%[[MAP0]], %[[MAP1]], %[[MAP2]] : !llvm.ptr, !llvm.ptr, !llvm.ptr)
+// CHECK: %[[MAP3:.*]] = omp.map_info var_ptr(%[[ARG0]] : !llvm.ptr, i32) map_clauses(from) capture(ByRef) -> !llvm.ptr {name = ""}
+// CHECK: %[[MAP4:.*]] = omp.map_info var_ptr(%[[ARG1]] : !llvm.ptr, i32) map_clauses(from) capture(ByRef) -> !llvm.ptr {name = ""}
+// CHECK: %[[MAP5:.*]] = omp.map_info var_ptr(%[[ARG2]] : !llvm.ptr, i32) map_clauses(exit_release_or_enter_alloc) capture(ByRef) -> !llvm.ptr {name = ""}
+// CHECK: %[[MAP6:.*]] = omp.map_info var_ptr(%[[ARG3]] : !llvm.ptr, i32) map_clauses(always, delete) capture(ByRef) -> !llvm.ptr {name = ""}
+// CHECK: omp.target_exit_data map_entries(%[[MAP3]], %[[MAP4]], %[[MAP5]], %[[MAP6]] : !llvm.ptr, !llvm.ptr, !llvm.ptr, !llvm.ptr)
+
+llvm.func @_QPomp_target_data(%a : !llvm.ptr, %b : !llvm.ptr, %c : !llvm.ptr, %d : !llvm.ptr) {
+ %0 = omp.map_info var_ptr(%a : !llvm.ptr, i32) map_clauses(to) capture(ByRef) -> !llvm.ptr {name = ""}
+ %1 = omp.map_info var_ptr(%b : !llvm.ptr, i32) map_clauses(to) capture(ByRef) -> !llvm.ptr {name = ""}
+ %2 = omp.map_info var_ptr(%c : !llvm.ptr, i32) map_clauses(always, exit_release_or_enter_alloc) capture(ByRef) -> !llvm.ptr {name = ""}
+ omp.target_enter_data map_entries(%0, %1, %2 : !llvm.ptr, !llvm.ptr, !llvm.ptr) {}
+ %3 = omp.map_info var_ptr(%a : !llvm.ptr, i32) map_clauses(from) capture(ByRef) -> !llvm.ptr {name = ""}
+ %4 = omp.map_info var_ptr(%b : !llvm.ptr, i32) map_clauses(from) capture(ByRef) -> !llvm.ptr {name = ""}
+ %5 = omp.map_info var_ptr(%c : !llvm.ptr, i32) map_clauses(exit_release_or_enter_alloc) capture(ByRef) -> !llvm.ptr {name = ""}
+ %6 = omp.map_info var_ptr(%d : !llvm.ptr, i32) map_clauses(always, delete) capture(ByRef) -> !llvm.ptr {name = ""}
+ omp.target_exit_data map_entries(%3, %4, %5, %6 : !llvm.ptr, !llvm.ptr, !llvm.ptr, !llvm.ptr) {}
llvm.return
}
// -----
// CHECK-LABEL: @_QPomp_target_data_region
-// CHECK: (%[[ARG0:.*]]: !llvm.ptr<array<1024 x i32>>, %[[ARG1:.*]]: !llvm.ptr<i32>) {
-// CHECK: %[[MAP_0:.*]] = omp.map_info var_ptr(%[[ARG0]] : !llvm.ptr<array<1024 x i32>>, !llvm.array<1024 x i32>) map_clauses(tofrom) capture(ByRef) -> !llvm.ptr<array<1024 x i32>> {name = ""}
-// CHECK: omp.target_data map_entries(%[[MAP_0]] : !llvm.ptr<array<1024 x i32>>) {
+// CHECK: (%[[ARG0:.*]]: !llvm.ptr, %[[ARG1:.*]]: !llvm.ptr) {
+// CHECK: %[[MAP_0:.*]] = omp.map_info var_ptr(%[[ARG0]] : !llvm.ptr, !llvm.array<1024 x i32>) map_clauses(tofrom) capture(ByRef) -> !llvm.ptr {name = ""}
+// CHECK: omp.target_data map_entries(%[[MAP_0]] : !llvm.ptr) {
// CHECK: %[[VAL_1:.*]] = llvm.mlir.constant(10 : i32) : i32
-// CHECK: llvm.store %[[VAL_1]], %[[ARG1]] : !llvm.ptr<i32>
+// CHECK: llvm.store %[[VAL_1]], %[[ARG1]] : i32, !llvm.ptr
// CHECK: omp.terminator
// CHECK: }
// CHECK: llvm.return
-llvm.func @_QPomp_target_data_region(%a : !llvm.ptr<array<1024 x i32>>, %i : !llvm.ptr<i32>) {
- %1 = omp.map_info var_ptr(%a : !llvm.ptr<array<1024 x i32>>, !llvm.array<1024 x i32>) map_clauses(tofrom) capture(ByRef) -> !llvm.ptr<array<1024 x i32>> {name = ""}
- omp.target_data map_entries(%1 : !llvm.ptr<array<1024 x i32>>) {
+llvm.func @_QPomp_target_data_region(%a : !llvm.ptr, %i : !llvm.ptr) {
+ %1 = omp.map_info var_ptr(%a : !llvm.ptr, !llvm.array<1024 x i32>) map_clauses(tofrom) capture(ByRef) -> !llvm.ptr {name = ""}
+ omp.target_data map_entries(%1 : !llvm.ptr) {
%2 = llvm.mlir.constant(10 : i32) : i32
- llvm.store %2, %i : !llvm.ptr<i32>
+ llvm.store %2, %i : i32, !llvm.ptr
omp.terminator
}
llvm.return
@@ -241,24 +241,24 @@ llvm.func @_QPomp_target_data_region(%a : !llvm.ptr<array<1024 x i32>>, %i : !ll
// -----
// CHECK-LABEL: llvm.func @_QPomp_target(
-// CHECK: %[[ARG_0:.*]]: !llvm.ptr<array<1024 x i32>>,
-// CHECK: %[[ARG_1:.*]]: !llvm.ptr<i32>) {
+// CHECK: %[[ARG_0:.*]]: !llvm.ptr,
+// CHECK: %[[ARG_1:.*]]: !llvm.ptr) {
// CHECK: %[[VAL_0:.*]] = llvm.mlir.constant(64 : i32) : i32
-// CHECK: %[[MAP:.*]] = omp.map_info var_ptr(%[[ARG_0]] : !llvm.ptr<array<1024 x i32>>, !llvm.array<1024 x i32>) map_clauses(tofrom) capture(ByRef) -> !llvm.ptr<array<1024 x i32>> {name = ""}
-// CHECK: omp.target thread_limit(%[[VAL_0]] : i32) map_entries(%[[MAP]] : !llvm.ptr<array<1024 x i32>>) {
+// CHECK: %[[MAP:.*]] = omp.map_info var_ptr(%[[ARG_0]] : !llvm.ptr, !llvm.array<1024 x i32>) map_clauses(tofrom) capture(ByRef) -> !llvm.ptr {name = ""}
+// CHECK: omp.target thread_limit(%[[VAL_0]] : i32) map_entries(%[[MAP]] : !llvm.ptr) {
// CHECK: %[[VAL_1:.*]] = llvm.mlir.constant(10 : i32) : i32
-// CHECK: llvm.store %[[VAL_1]], %[[ARG_1]] : !llvm.ptr<i32>
+// CHECK: llvm.store %[[VAL_1]], %[[ARG_1]] : i32, !llvm.ptr
// CHECK: omp.terminator
// CHECK: }
// CHECK: llvm.return
// CHECK: }
-llvm.func @_QPomp_target(%a : !llvm.ptr<array<1024 x i32>>, %i : !llvm.ptr<i32>) {
+llvm.func @_QPomp_target(%a : !llvm.ptr, %i : !llvm.ptr) {
%0 = llvm.mlir.constant(64 : i32) : i32
- %1 = omp.map_info var_ptr(%a : !llvm.ptr<array<1024 x i32>>, !llvm.array<1024 x i32>) map_clauses(tofrom) capture(ByRef) -> !llvm.ptr<array<1024 x i32>> {name = ""}
- omp.target thread_limit(%0 : i32) map_entries(%1 : !llvm.ptr<array<1024 x i32>>) {
+ %1 = omp.map_info var_ptr(%a : !llvm.ptr, !llvm.array<1024 x i32>) map_clauses(tofrom) capture(ByRef) -> !llvm.ptr {name = ""}
+ omp.target thread_limit(%0 : i32) map_entries(%1 : !llvm.ptr) {
%2 = llvm.mlir.constant(10 : i32) : i32
- llvm.store %2, %i : !llvm.ptr<i32>
+ llvm.store %2, %i : i32, !llvm.ptr
omp.terminator
}
llvm.return
@@ -314,10 +314,10 @@ llvm.func @_QPsb() {
// CHECK: %[[COMBINE_VAL_EXT:.*]] = llvm.zext %[[COMBINE_VAL]] : i1 to i32
// CHECK: omp.yield(%[[COMBINE_VAL_EXT]] : i32)
// CHECK-LABEL: @_QPsimple_reduction
-// CHECK: %[[RED_ACCUMULATOR:.*]] = llvm.alloca %{{.*}} x i32 {bindc_name = "x", uniq_name = "_QFsimple_reductionEx"} : (i64) -> !llvm.ptr<i32>
+// CHECK: %[[RED_ACCUMULATOR:.*]] = llvm.alloca %{{.*}} x i32 {bindc_name = "x", uniq_name = "_QFsimple_reductionEx"} : (i64) -> !llvm.ptr
// CHECK: omp.parallel
-// CHECK: omp.wsloop reduction(@eqv_reduction -> %[[RED_ACCUMULATOR]] : !llvm.ptr<i32>) for
-// CHECK: omp.reduction %{{.*}}, %[[RED_ACCUMULATOR]] : i32, !llvm.ptr<i32>
+// CHECK: omp.wsloop reduction(@eqv_reduction -> %[[RED_ACCUMULATOR]] : !llvm.ptr) for
+// CHECK: omp.reduction %{{.*}}, %[[RED_ACCUMULATOR]] : i32, !llvm.ptr
// CHECK: omp.yield
// CHECK: omp.terminator
// CHECK: llvm.return
@@ -336,24 +336,24 @@ omp.reduction.declare @eqv_reduction : i32 init {
%4 = llvm.zext %3 : i1 to i32
omp.yield(%4 : i32)
}
-llvm.func @_QPsimple_reduction(%arg0: !llvm.ptr<array<100 x i32>> {fir.bindc_name = "y"}) {
+llvm.func @_QPsimple_reduction(%arg0: !llvm.ptr {fir.bindc_name = "y"}) {
%0 = llvm.mlir.constant(100 : i32) : i32
%1 = llvm.mlir.constant(1 : i32) : i32
%2 = llvm.mlir.constant(true) : i1
%3 = llvm.mlir.constant(1 : i64) : i64
- %4 = llvm.alloca %3 x i32 {bindc_name = "x", uniq_name = "_QFsimple_reductionEx"} : (i64) -> !llvm.ptr<i32>
+ %4 = llvm.alloca %3 x i32 {bindc_name = "x", uniq_name = "_QFsimple_reductionEx"} : (i64) -> !llvm.ptr
%5 = llvm.zext %2 : i1 to i32
- llvm.store %5, %4 : !llvm.ptr<i32>
+ llvm.store %5, %4 : i32, !llvm.ptr
omp.parallel {
- %6 = llvm.alloca %3 x i32 {adapt.valuebyref, in_type = i32, operandSegmentSizes = array<i32: 0, 0>, pinned} : (i64) -> !llvm.ptr<i32>
- omp.wsloop reduction(@eqv_reduction -> %4 : !llvm.ptr<i32>) for (%arg1) : i32 = (%1) to (%0) inclusive step (%1) {
- llvm.store %arg1, %6 : !llvm.ptr<i32>
- %7 = llvm.load %6 : !llvm.ptr<i32>
+ %6 = llvm.alloca %3 x i32 {adapt.valuebyref, in_type = i32, operandSegmentSizes = array<i32: 0, 0>, pinned} : (i64) -> !llvm.ptr
+ omp.wsloop reduction(@eqv_reduction -> %4 : !llvm.ptr) for (%arg1) : i32 = (%1) to (%0) inclusive step (%1) {
+ llvm.store %arg1, %6 : i32, !llvm.ptr
+ %7 = llvm.load %6 : !llvm.ptr -> i32
%8 = llvm.sext %7 : i32 to i64
%9 = llvm.sub %8, %3 : i64
- %10 = llvm.getelementptr %arg0[0, %9] : (!llvm.ptr<array<100 x i32>>, i64) -> !llvm.ptr<i32>
- %11 = llvm.load %10 : !llvm.ptr<i32>
- omp.reduction %11, %4 : i32, !llvm.ptr<i32>
+ %10 = llvm.getelementptr %arg0[0, %9] : (!llvm.ptr, i64) -> !llvm.ptr, !llvm.array<100 x i32>
+ %11 = llvm.load %10 : !llvm.ptr -> i32
+ omp.reduction %11, %4 : i32, !llvm.ptr
omp.yield
}
omp.terminator
@@ -369,7 +369,7 @@ llvm.func @_QQmain() {
%1 = llvm.mlir.constant(5 : index) : i64
%2 = llvm.mlir.constant(1 : index) : i64
%3 = llvm.mlir.constant(1 : i64) : i64
- %4 = llvm.alloca %3 x i32 : (i64) -> !llvm.ptr<i32>
+ %4 = llvm.alloca %3 x i32 : (i64) -> !llvm.ptr
// CHECK: omp.taskgroup
omp.taskgroup {
%5 = llvm.trunc %2 : i64 to i32
@@ -378,27 +378,27 @@ llvm.func @_QQmain() {
%8 = llvm.icmp "sgt" %7, %0 : i64
llvm.cond_br %8, ^bb2, ^bb3
^bb2: // pred: ^bb1
- llvm.store %6, %4 : !llvm.ptr<i32>
+ llvm.store %6, %4 : i32, !llvm.ptr
// CHECK: omp.task
omp.task {
// CHECK: llvm.call @[[CALL_FUNC:.*]]({{.*}}) :
- llvm.call @_QFPdo_work(%4) : (!llvm.ptr<i32>) -> ()
+ llvm.call @_QFPdo_work(%4) : (!llvm.ptr) -> ()
// CHECK: omp.terminator
omp.terminator
}
- %9 = llvm.load %4 : !llvm.ptr<i32>
+ %9 = llvm.load %4 : !llvm.ptr -> i32
%10 = llvm.add %9, %5 : i32
%11 = llvm.sub %7, %2 : i64
llvm.br ^bb1(%10, %11 : i32, i64)
^bb3: // pred: ^bb1
- llvm.store %6, %4 : !llvm.ptr<i32>
+ llvm.store %6, %4 : i32, !llvm.ptr
// CHECK: omp.terminator
omp.terminator
}
llvm.return
}
// CHECK: @[[CALL_FUNC]]
-llvm.func @_QFPdo_work(%arg0: !llvm.ptr<i32> {fir.bindc_name = "i"}) {
+llvm.func @_QFPdo_work(%arg0: !llvm.ptr {fir.bindc_name = "i"}) {
llvm.return
}
@@ -409,7 +409,7 @@ llvm.func @sub_() {
%0 = llvm.mlir.constant(0 : index) : i64
%1 = llvm.mlir.constant(1 : index) : i64
%2 = llvm.mlir.constant(1 : i64) : i64
- %3 = llvm.alloca %2 x i32 {bindc_name = "i", in_type = i32, operandSegmentSizes = array<i32: 0, 0>, uniq_name = "_QFsubEi"} : (i64) -> !llvm.ptr<i32>
+ %3 = llvm.alloca %2 x i32 {bindc_name = "i", in_type = i32, operandSegmentSizes = array<i32: 0, 0>, uniq_name = "_QFsubEi"} : (i64) -> !llvm.ptr
// CHECK: omp.ordered_region
omp.ordered_region {
%4 = llvm.trunc %1 : i64 to i32
@@ -418,15 +418,15 @@ llvm.func @sub_() {
%7 = llvm.icmp "sgt" %6, %0 : i64
llvm.cond_br %7, ^bb2, ^bb3
^bb2: // pred: ^bb1
- llvm.store %5, %3 : !llvm.ptr<i32>
- %8 = llvm.load %3 : !llvm.ptr<i32>
+ llvm.store %5, %3 : i32, !llvm.ptr
+ %8 = llvm.load %3 : !llvm.ptr -> i32
// CHECK: llvm.add
%9 = arith.addi %8, %4 : i32
// CHECK: llvm.sub
%10 = arith.subi %6, %1 : i64
llvm.br ^bb1(%9, %10 : i32, i64)
^bb3: // pred: ^bb1
- llvm.store %5, %3 : !llvm.ptr<i32>
+ llvm.store %5, %3 : i32, !llvm.ptr
// CHECK: omp.terminator
omp.terminator
}
@@ -436,41 +436,39 @@ llvm.func @sub_() {
// -----
// CHECK-LABEL: llvm.func @_QPtarget_map_with_bounds(
-// CHECK: %[[ARG_0:.*]]: !llvm.ptr<i32>,
-// CHECK: %[[ARG_1:.*]]: !llvm.ptr<array<10 x i32>>,
-// CHECK: %[[ARG_2:.*]]: !llvm.ptr<array<10 x i32>>) {
+// CHECK: %[[ARG_0:.*]]: !llvm.ptr, %[[ARG_1:.*]]: !llvm.ptr, %[[ARG_2:.*]]: !llvm.ptr) {
// CHECK: %[[C_01:.*]] = llvm.mlir.constant(4 : index) : i64
// CHECK: %[[C_02:.*]] = llvm.mlir.constant(1 : index) : i64
// CHECK: %[[C_03:.*]] = llvm.mlir.constant(1 : index) : i64
// CHECK: %[[C_04:.*]] = llvm.mlir.constant(1 : index) : i64
// CHECK: %[[BOUNDS0:.*]] = omp.bounds lower_bound(%[[C_02]] : i64) upper_bound(%[[C_01]] : i64) stride(%[[C_04]] : i64) start_idx(%[[C_04]] : i64)
-// CHECK: %[[MAP0:.*]] = omp.map_info var_ptr(%[[ARG_1]] : !llvm.ptr<array<10 x i32>>, !llvm.array<10 x i32>) map_clauses(tofrom) capture(ByRef) bounds(%[[BOUNDS0]]) -> !llvm.ptr<array<10 x i32>> {name = ""}
+// CHECK: %[[MAP0:.*]] = omp.map_info var_ptr(%[[ARG_1]] : !llvm.ptr, !llvm.array<10 x i32>) map_clauses(tofrom) capture(ByRef) bounds(%[[BOUNDS0]]) -> !llvm.ptr {name = ""}
// CHECK: %[[C_11:.*]] = llvm.mlir.constant(4 : index) : i64
// CHECK: %[[C_12:.*]] = llvm.mlir.constant(1 : index) : i64
// CHECK: %[[C_13:.*]] = llvm.mlir.constant(1 : index) : i64
// CHECK: %[[C_14:.*]] = llvm.mlir.constant(1 : index) : i64
// CHECK: %[[BOUNDS1:.*]] = omp.bounds lower_bound(%[[C_12]] : i64) upper_bound(%[[C_11]] : i64) stride(%[[C_14]] : i64) start_idx(%[[C_14]] : i64)
-// CHECK: %[[MAP1:.*]] = omp.map_info var_ptr(%[[ARG_2]] : !llvm.ptr<array<10 x i32>>, !llvm.array<10 x i32>) map_clauses(tofrom) capture(ByRef) bounds(%[[BOUNDS1]]) -> !llvm.ptr<array<10 x i32>> {name = ""}
-// CHECK: omp.target map_entries(%[[MAP0]], %[[MAP1]] : !llvm.ptr<array<10 x i32>>, !llvm.ptr<array<10 x i32>>) {
+// CHECK: %[[MAP1:.*]] = omp.map_info var_ptr(%[[ARG_2]] : !llvm.ptr, !llvm.array<10 x i32>) map_clauses(tofrom) capture(ByRef) bounds(%[[BOUNDS1]]) -> !llvm.ptr {name = ""}
+// CHECK: omp.target map_entries(%[[MAP0]], %[[MAP1]] : !llvm.ptr, !llvm.ptr) {
// CHECK: omp.terminator
// CHECK: }
// CHECK: llvm.return
// CHECK:}
-llvm.func @_QPtarget_map_with_bounds(%arg0: !llvm.ptr<i32>, %arg1: !llvm.ptr<array<10 x i32>>, %arg2: !llvm.ptr<array<10 x i32>>) {
+llvm.func @_QPtarget_map_with_bounds(%arg0: !llvm.ptr, %arg1: !llvm.ptr, %arg2: !llvm.ptr) {
%0 = llvm.mlir.constant(4 : index) : i64
%1 = llvm.mlir.constant(1 : index) : i64
%2 = llvm.mlir.constant(1 : index) : i64
%3 = llvm.mlir.constant(1 : index) : i64
%4 = omp.bounds lower_bound(%1 : i64) upper_bound(%0 : i64) stride(%3 : i64) start_idx(%3 : i64)
- %5 = omp.map_info var_ptr(%arg1 : !llvm.ptr<array<10 x i32>>, !llvm.array<10 x i32>) map_clauses(tofrom) capture(ByRef) bounds(%4) -> !llvm.ptr<array<10 x i32>> {name = ""}
+ %5 = omp.map_info var_ptr(%arg1 : !llvm.ptr, !llvm.array<10 x i32>) map_clauses(tofrom) capture(ByRef) bounds(%4) -> !llvm.ptr {name = ""}
%6 = llvm.mlir.constant(4 : index) : i64
%7 = llvm.mlir.constant(1 : index) : i64
%8 = llvm.mlir.constant(1 : index) : i64
%9 = llvm.mlir.constant(1 : index) : i64
%10 = omp.bounds lower_bound(%7 : i64) upper_bound(%6 : i64) stride(%9 : i64) start_idx(%9 : i64)
- %11 = omp.map_info var_ptr(%arg2 : !llvm.ptr<array<10 x i32>>, !llvm.array<10 x i32>) map_clauses(tofrom) capture(ByRef) bounds(%10) -> !llvm.ptr<array<10 x i32>> {name = ""}
- omp.target map_entries(%5, %11 : !llvm.ptr<array<10 x i32>>, !llvm.ptr<array<10 x i32>>) {
+ %11 = omp.map_info var_ptr(%arg2 : !llvm.ptr, !llvm.array<10 x i32>) map_clauses(tofrom) capture(ByRef) bounds(%10) -> !llvm.ptr {name = ""}
+ omp.target map_entries(%5, %11 : !llvm.ptr, !llvm.ptr) {
omp.terminator
}
llvm.return
diff --git a/mlir/test/Dialect/GPU/invalid.mlir b/mlir/test/Dialect/GPU/invalid.mlir
index c8c0b7d24bc3ab2..680e604151d77fd 100644
--- a/mlir/test/Dialect/GPU/invalid.mlir
+++ b/mlir/test/Dialect/GPU/invalid.mlir
@@ -88,7 +88,7 @@ module attributes {gpu.container_module} {
module attributes {gpu.container_module} {
module @kernels {
// expected-error at +1 {{'gpu.func' op expects parent op 'gpu.module'}}
- gpu.func @kernel_1(%arg1 : !llvm.ptr<f32>) {
+ gpu.func @kernel_1(%arg1 : !llvm.ptr) {
gpu.return
}
}
@@ -138,14 +138,14 @@ module attributes {gpu.container_module} {
module attributes {gpu.container_module} {
module @kernels {
- gpu.func @kernel_1(%arg1 : !llvm.ptr<f32>) kernel {
+ gpu.func @kernel_1(%arg1 : !llvm.ptr) kernel {
gpu.return
}
}
- func.func @launch_func_missing_kernel_attr(%sz : index, %arg : !llvm.ptr<f32>) {
+ func.func @launch_func_missing_kernel_attr(%sz : index, %arg : !llvm.ptr) {
// expected-error at +1 {{kernel module 'kernels' is undefined}}
- gpu.launch_func @kernels::@kernel_1 blocks in (%sz, %sz, %sz) threads in (%sz, %sz, %sz) args(%arg : !llvm.ptr<f32>)
+ gpu.launch_func @kernels::@kernel_1 blocks in (%sz, %sz, %sz) threads in (%sz, %sz, %sz) args(%arg : !llvm.ptr)
return
}
}
@@ -154,14 +154,14 @@ module attributes {gpu.container_module} {
module attributes {gpu.container_module} {
gpu.module @kernels {
- gpu.func @kernel_1(%arg1 : !llvm.ptr<f32>) {
+ gpu.func @kernel_1(%arg1 : !llvm.ptr) {
gpu.return
}
}
- func.func @launch_func_missing_kernel_attr(%sz : index, %arg : !llvm.ptr<f32>) {
+ func.func @launch_func_missing_kernel_attr(%sz : index, %arg : !llvm.ptr) {
// expected-error at +1 {{kernel function is missing the 'gpu.kernel' attribute}}
- gpu.launch_func @kernels::@kernel_1 blocks in (%sz, %sz, %sz) threads in (%sz, %sz, %sz) args(%arg : !llvm.ptr<f32>)
+ gpu.launch_func @kernels::@kernel_1 blocks in (%sz, %sz, %sz) threads in (%sz, %sz, %sz) args(%arg : !llvm.ptr)
return
}
}
@@ -170,14 +170,14 @@ module attributes {gpu.container_module} {
module attributes {gpu.container_module} {
gpu.module @kernels {
- gpu.func @kernel_1(%arg1 : !llvm.ptr<f32>) kernel {
+ gpu.func @kernel_1(%arg1 : !llvm.ptr) kernel {
gpu.return
}
}
- func.func @launch_func_kernel_operand_size(%sz : index, %arg : !llvm.ptr<f32>) {
+ func.func @launch_func_kernel_operand_size(%sz : index, %arg : !llvm.ptr) {
// expected-error at +1 {{got 2 kernel operands but expected 1}}
- gpu.launch_func @kernels::@kernel_1 blocks in (%sz, %sz, %sz) threads in (%sz, %sz, %sz) args(%arg : !llvm.ptr<f32>, %arg : !llvm.ptr<f32>)
+ gpu.launch_func @kernels::@kernel_1 blocks in (%sz, %sz, %sz) threads in (%sz, %sz, %sz) args(%arg : !llvm.ptr, %arg : !llvm.ptr)
return
}
}
diff --git a/mlir/test/Dialect/GPU/module-to-binary-nvvm.mlir b/mlir/test/Dialect/GPU/module-to-binary-nvvm.mlir
index 22d7caa38feec97..05e368f7a642e66 100644
--- a/mlir/test/Dialect/GPU/module-to-binary-nvvm.mlir
+++ b/mlir/test/Dialect/GPU/module-to-binary-nvvm.mlir
@@ -6,8 +6,8 @@ module attributes {gpu.container_module} {
// CHECK-LABEL:gpu.binary @kernel_module1
// CHECK:[#gpu.object<#nvvm.target<chip = "sm_70">, offload = "{{.*}}">]
gpu.module @kernel_module1 [#nvvm.target<chip = "sm_70">] {
- llvm.func @kernel(%arg0: i32, %arg1: !llvm.ptr<f32>,
- %arg2: !llvm.ptr<f32>, %arg3: i64, %arg4: i64,
+ llvm.func @kernel(%arg0: i32, %arg1: !llvm.ptr,
+ %arg2: !llvm.ptr, %arg3: i64, %arg4: i64,
%arg5: i64) attributes {gpu.kernel} {
llvm.return
}
@@ -16,8 +16,8 @@ module attributes {gpu.container_module} {
// CHECK-LABEL:gpu.binary @kernel_module2
// CHECK-ISA:[#gpu.object<#nvvm.target<flags = {fast}>, properties = {O = 2 : i32}, assembly = "{{.*}}">, #gpu.object<#nvvm.target, properties = {O = 2 : i32}, assembly = "{{.*}}">]
gpu.module @kernel_module2 [#nvvm.target<flags = {fast}>, #nvvm.target] {
- llvm.func @kernel(%arg0: i32, %arg1: !llvm.ptr<f32>,
- %arg2: !llvm.ptr<f32>, %arg3: i64, %arg4: i64,
+ llvm.func @kernel(%arg0: i32, %arg1: !llvm.ptr,
+ %arg2: !llvm.ptr, %arg3: i64, %arg4: i64,
%arg5: i64) attributes {gpu.kernel} {
llvm.return
}
diff --git a/mlir/test/Dialect/GPU/module-to-binary-rocdl.mlir b/mlir/test/Dialect/GPU/module-to-binary-rocdl.mlir
index 9f987c71387f4ca..939dbdd4382e73b 100644
--- a/mlir/test/Dialect/GPU/module-to-binary-rocdl.mlir
+++ b/mlir/test/Dialect/GPU/module-to-binary-rocdl.mlir
@@ -6,8 +6,8 @@ module attributes {gpu.container_module} {
// CHECK-LABEL:gpu.binary @kernel_module1
// CHECK:[#gpu.object<#rocdl.target<chip = "gfx90a">, offload = "{{.*}}">]
gpu.module @kernel_module1 [#rocdl.target<chip = "gfx90a">] {
- llvm.func @kernel(%arg0: i32, %arg1: !llvm.ptr<f32>,
- %arg2: !llvm.ptr<f32>, %arg3: i64, %arg4: i64,
+ llvm.func @kernel(%arg0: i32, %arg1: !llvm.ptr,
+ %arg2: !llvm.ptr, %arg3: i64, %arg4: i64,
%arg5: i64) attributes {gpu.kernel} {
llvm.return
}
@@ -16,8 +16,8 @@ module attributes {gpu.container_module} {
// CHECK-LABEL:gpu.binary @kernel_module2
// CHECK-ISA:[#gpu.object<#rocdl.target<flags = {fast}>, assembly = "{{.*}}">, #gpu.object<#rocdl.target, assembly = "{{.*}}">]
gpu.module @kernel_module2 [#rocdl.target<flags = {fast}>, #rocdl.target] {
- llvm.func @kernel(%arg0: i32, %arg1: !llvm.ptr<f32>,
- %arg2: !llvm.ptr<f32>, %arg3: i64, %arg4: i64,
+ llvm.func @kernel(%arg0: i32, %arg1: !llvm.ptr,
+ %arg2: !llvm.ptr, %arg3: i64, %arg4: i64,
%arg5: i64) attributes {gpu.kernel} {
llvm.return
}
diff --git a/mlir/test/Dialect/GPU/outlining.mlir b/mlir/test/Dialect/GPU/outlining.mlir
index ca776968d998b41..28c121a550100c2 100644
--- a/mlir/test/Dialect/GPU/outlining.mlir
+++ b/mlir/test/Dialect/GPU/outlining.mlir
@@ -263,7 +263,7 @@ func.func @function_call(%arg0 : memref<?xf32>) {
%block_z = %cst) {
func.call @device_function() : () -> ()
func.call @device_function() : () -> ()
- %0 = llvm.mlir.addressof @global : !llvm.ptr<i64>
+ %0 = llvm.mlir.addressof @global : !llvm.ptr
gpu.terminator
}
return
@@ -285,7 +285,7 @@ func.func @recursive_device_function() {
// CHECK: gpu.func @function_call_kernel()
// CHECK: call @device_function() : () -> ()
// CHECK: call @device_function() : () -> ()
-// CHECK: llvm.mlir.addressof @global : !llvm.ptr<i64>
+// CHECK: llvm.mlir.addressof @global : !llvm.ptr
// CHECK: gpu.return
//
// CHECK: llvm.mlir.global internal @global(42 : i64) {addr_space = 0 : i32} : i64
diff --git a/mlir/test/Dialect/LLVMIR/canonicalize-typed-pointers.mlir b/mlir/test/Dialect/LLVMIR/canonicalize-typed-pointers.mlir
deleted file mode 100644
index 2ae9727482fc331..000000000000000
--- a/mlir/test/Dialect/LLVMIR/canonicalize-typed-pointers.mlir
+++ /dev/null
@@ -1,86 +0,0 @@
-// 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/dynamic-gep-index-typed-pointers.mlir b/mlir/test/Dialect/LLVMIR/dynamic-gep-index-typed-pointers.mlir
deleted file mode 100644
index 9e14b1db3432b0e..000000000000000
--- a/mlir/test/Dialect/LLVMIR/dynamic-gep-index-typed-pointers.mlir
+++ /dev/null
@@ -1,12 +0,0 @@
-// 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/global-typed-pointers.mlir b/mlir/test/Dialect/LLVMIR/global-typed-pointers.mlir
deleted file mode 100644
index 56d720cc866b6f7..000000000000000
--- a/mlir/test/Dialect/LLVMIR/global-typed-pointers.mlir
+++ /dev/null
@@ -1,46 +0,0 @@
-// 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/invalid-typed-pointers.mlir b/mlir/test/Dialect/LLVMIR/invalid-typed-pointers.mlir
deleted file mode 100644
index a87b1952b6dca70..000000000000000
--- a/mlir/test/Dialect/LLVMIR/invalid-typed-pointers.mlir
+++ /dev/null
@@ -1,299 +0,0 @@
-// 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 @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
-
-// expected-error at below {{'llvm.resume' should have a consistent input type inside a function}}
-llvm.func @caller(%arg0: i32, %arg1: !llvm.struct<(ptr<i32>, i32)>) -> i32 attributes { personality = @__gxx_personality_v0 } {
- %0 = llvm.invoke @foo(%arg0) to ^bb1 unwind ^bb2 : (i32) -> i32
-^bb1:
- %1 = llvm.invoke @foo(%0) to ^bb3 unwind ^bb4 : (i32) -> i32
-^bb2:
- %2 = llvm.landingpad cleanup : !llvm.struct<(ptr<i8>, i32)>
- llvm.resume %arg1 : !llvm.struct<(ptr<i32>, i32)>
-^bb3:
- llvm.return %1 : i32
-^bb4:
- %3 = llvm.landingpad cleanup : !llvm.struct<(ptr<i8>, i32)>
- llvm.resume %3 : !llvm.struct<(ptr<i8>, i32)>
-}
-
-// -----
-
-llvm.func @foo(i32) -> i32
-llvm.func @__gxx_personality_v0(...) -> i32
-
-// expected-error at below {{'llvm.landingpad' should have a consistent result type inside a function}}
-llvm.func @caller(%arg0: i32) -> i32 attributes { personality = @__gxx_personality_v0 } {
- %0 = llvm.invoke @foo(%arg0) to ^bb1 unwind ^bb2 : (i32) -> i32
-^bb1:
- %1 = llvm.invoke @foo(%0) to ^bb3 unwind ^bb4 : (i32) -> i32
-^bb2:
- %2 = llvm.landingpad cleanup : !llvm.struct<(ptr<i8>, i32)>
- llvm.resume %2 : !llvm.struct<(ptr<i8>, i32)>
-^bb3:
- llvm.return %1 : i32
-^bb4:
- %3 = llvm.landingpad cleanup : !llvm.struct<(ptr<i32>, i32)>
- llvm.resume %3 : !llvm.struct<(ptr<i32>, 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, cache = ca : !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 {{CG cache modifier is only support for 16 bytes copy.}}
- nvvm.cp.async.shared.global %arg0, %arg1, 8, cache = cg : !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/layout-typed-pointers.mlir b/mlir/test/Dialect/LLVMIR/layout-typed-pointers.mlir
deleted file mode 100644
index 5cf1ed03e64c824..000000000000000
--- a/mlir/test/Dialect/LLVMIR/layout-typed-pointers.mlir
+++ /dev/null
@@ -1,145 +0,0 @@
-// 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/nvvm-typed-pointers.mlir b/mlir/test/Dialect/LLVMIR/nvvm-typed-pointers.mlir
deleted file mode 100644
index 68eb1ecca00b884..000000000000000
--- a/mlir/test/Dialect/LLVMIR/nvvm-typed-pointers.mlir
+++ /dev/null
@@ -1,55 +0,0 @@
-// 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 %{{.*}}, %{{.*}}, cache = ca
- nvvm.cp.async.shared.global %arg0, %arg1, 16, cache=ca : !llvm.ptr<i8, 3>, !llvm.ptr<i8, 1>
-// CHECK: nvvm.cp.async.shared.global %{{.*}}, %{{.*}}, cache = cg
- nvvm.cp.async.shared.global %arg0, %arg1, 16, cache=cg : !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/parameter-attrs-invalid-typed-pointers.mlir b/mlir/test/Dialect/LLVMIR/parameter-attrs-invalid-typed-pointers.mlir
deleted file mode 100644
index 65411ff41e28515..000000000000000
--- a/mlir/test/Dialect/LLVMIR/parameter-attrs-invalid-typed-pointers.mlir
+++ /dev/null
@@ -1,6 +0,0 @@
-// RUN: mlir-opt %s -split-input-file -verify-diagnostics
-
-// Argument attributes
-
-// expected-error at below {{"llvm.sret" attribute attached to LLVM pointer argument of different 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 58b48345528411e..55b1d4faf207f96 100644
--- a/mlir/test/Dialect/LLVMIR/parameter-attrs-invalid.mlir
+++ b/mlir/test/Dialect/LLVMIR/parameter-attrs-invalid.mlir
@@ -52,31 +52,16 @@ llvm.func @invalid_byval_arg_type(%0 : i32 {llvm.byval = !llvm.struct<(i32)>})
// -----
-// expected-error at below {{"llvm.byval" attribute attached to LLVM pointer argument of different type}}
-llvm.func @invalid_byval_attr_type(%0 : !llvm.ptr<!llvm.struct<(f32)>> {llvm.byval = !llvm.struct<(i32)>})
-
-// -----
-
// expected-error at below {{"llvm.byref" attribute attached to non-pointer LLVM type}}
llvm.func @invalid_byref_arg_type(%0 : i32 {llvm.byref = !llvm.struct<(i32)>})
// -----
-// expected-error at below {{"llvm.byref" attribute attached to LLVM pointer argument of different type}}
-llvm.func @invalid_byref_attr_type(%0 : !llvm.ptr<!llvm.struct<(f32)>> {llvm.byref = !llvm.struct<(i32)>})
-
-// -----
-
// expected-error at below {{"llvm.inalloca" attribute attached to non-pointer LLVM type}}
llvm.func @invalid_inalloca_arg_type(%0 : i32 {llvm.inalloca = !llvm.struct<(i32)>})
// -----
-// expected-error at below {{"llvm.inalloca" attribute attached to LLVM pointer argument of different type}}
-llvm.func @invalid_inalloca_attr_type(%0 : !llvm.ptr<!llvm.struct<(f32)>> {llvm.inalloca = !llvm.struct<(i32)>})
-
-// -----
-
// expected-error at below {{"llvm.signext" attribute attached to non-integer LLVM type}}
llvm.func @invalid_signext_arg_type(%0 : f32 {llvm.signext})
diff --git a/mlir/test/Dialect/LLVMIR/roundtrip-typed-pointers.mlir b/mlir/test/Dialect/LLVMIR/roundtrip-typed-pointers.mlir
deleted file mode 100644
index f974bcd2e02aff6..000000000000000
--- a/mlir/test/Dialect/LLVMIR/roundtrip-typed-pointers.mlir
+++ /dev/null
@@ -1,73 +0,0 @@
-// RUN: mlir-opt %s | mlir-opt | FileCheck %s
-
-// CHECK-LABEL: func @ops
-// CHECK-SAME: %[[I32:.*]]: i32
-func.func @ops(%arg0: i32) {
-// Memory-related operations.
-//
-// CHECK-NEXT: %[[ALLOCA:.*]] = llvm.alloca %[[I32]] x f64 : (i32) -> !llvm.ptr<f64>
-// CHECK-NEXT: %[[GEP:.*]] = llvm.getelementptr %[[ALLOCA]][%[[I32]]] : (!llvm.ptr<f64>, i32) -> !llvm.ptr<f64>
-// CHECK-NEXT: %[[VALUE:.*]] = llvm.load %[[GEP]] : !llvm.ptr<f64>
-// CHECK-NEXT: llvm.store %[[VALUE]], %[[ALLOCA]] : !llvm.ptr<f64>
-// CHECK-NEXT: %{{.*}} = llvm.bitcast %[[ALLOCA]] : !llvm.ptr<f64> to !llvm.ptr<i64>
- %13 = llvm.alloca %arg0 x f64 : (i32) -> !llvm.ptr<f64>
- %14 = llvm.getelementptr %13[%arg0] : (!llvm.ptr<f64>, i32) -> !llvm.ptr<f64>
- %15 = llvm.load %14 : !llvm.ptr<f64>
- llvm.store %15, %13 : !llvm.ptr<f64>
- %16 = llvm.bitcast %13 : !llvm.ptr<f64> to !llvm.ptr<i64>
- llvm.return
-}
-
-// CHECK-LABEL: @gep
-llvm.func @gep(%ptr: !llvm.ptr<struct<(i32, struct<(i32, f32)>)>>, %idx: i64,
- %ptr2: !llvm.ptr<struct<(array<10 x f32>)>>) {
- // CHECK: llvm.getelementptr %{{.*}}[%{{.*}}, 1, 0] : (!llvm.ptr<struct<(i32, struct<(i32, f32)>)>>, i64) -> !llvm.ptr<i32>
- llvm.getelementptr %ptr[%idx, 1, 0] : (!llvm.ptr<struct<(i32, struct<(i32, f32)>)>>, i64) -> !llvm.ptr<i32>
- // CHECK: llvm.getelementptr inbounds %{{.*}}[%{{.*}}, 0, %{{.*}}] : (!llvm.ptr<struct<(array<10 x f32>)>>, i64, i64) -> !llvm.ptr<f32>
- llvm.getelementptr inbounds %ptr2[%idx, 0, %idx] : (!llvm.ptr<struct<(array<10 x f32>)>>, i64, i64) -> !llvm.ptr<f32>
- llvm.return
-}
-
-// CHECK-LABEL: @alloca
-func.func @alloca(%size : i64) {
- // CHECK: llvm.alloca %{{.*}} x i32 : (i64) -> !llvm.ptr<i32>
- llvm.alloca %size x i32 {alignment = 0} : (i64) -> (!llvm.ptr<i32>)
- // CHECK: llvm.alloca inalloca %{{.*}} x i32 {alignment = 8 : i64} : (i64) -> !llvm.ptr<i32>
- llvm.alloca inalloca %size x i32 {alignment = 8} : (i64) -> (!llvm.ptr<i32>)
- llvm.return
-}
-
-// CHECK-LABEL: @null
-func.func @null() {
- // CHECK: llvm.mlir.zero : !llvm.ptr<i8>
- %0 = llvm.mlir.zero : !llvm.ptr<i8>
- // CHECK: llvm.mlir.zero : !llvm.ptr<struct<(ptr<func<void (i32, ptr<func<void ()>>)>>, i64)>>
- %1 = llvm.mlir.zero : !llvm.ptr<struct<(ptr<func<void (i32, ptr<func<void ()>>)>>, i64)>>
- llvm.return
-}
-
-// CHECK-LABEL: llvm.func @vararg_func
-llvm.func @vararg_func(%arg0: i32, ...) {
- // CHECK: %{{.*}} = llvm.mlir.constant(1 : i32) : i32
- // CHECK: %{{.*}} = llvm.mlir.constant(1 : i32) : i32
- %0 = llvm.mlir.constant(1 : i32) : i32
- %1 = llvm.mlir.constant(1 : i32) : i32
- // CHECK: %[[ALLOCA0:.+]] = llvm.alloca %{{.*}} x !llvm.struct<"struct.va_list", (ptr<i8>)> {alignment = 8 : i64} : (i32) -> !llvm.ptr<struct<"struct.va_list", (ptr<i8>)>>
- // CHECK: %[[CAST0:.+]] = llvm.bitcast %[[ALLOCA0]] : !llvm.ptr<struct<"struct.va_list", (ptr<i8>)>> to !llvm.ptr<i8>
- %2 = llvm.alloca %1 x !llvm.struct<"struct.va_list", (ptr<i8>)> {alignment = 8 : i64} : (i32) -> !llvm.ptr<struct<"struct.va_list", (ptr<i8>)>>
- %3 = llvm.bitcast %2 : !llvm.ptr<struct<"struct.va_list", (ptr<i8>)>> to !llvm.ptr<i8>
- // CHECK: llvm.intr.vastart %[[CAST0]]
- llvm.intr.vastart %3 : !llvm.ptr<i8>
- // CHECK: %[[ALLOCA1:.+]] = llvm.alloca %{{.*}} x !llvm.ptr<i8> {alignment = 8 : i64} : (i32) -> !llvm.ptr<ptr<i8>>
- // CHECK: %[[CAST1:.+]] = llvm.bitcast %[[ALLOCA1]] : !llvm.ptr<ptr<i8>> to !llvm.ptr<i8>
- %4 = llvm.alloca %0 x !llvm.ptr<i8> {alignment = 8 : i64} : (i32) -> !llvm.ptr<ptr<i8>>
- %5 = llvm.bitcast %4 : !llvm.ptr<ptr<i8>> to !llvm.ptr<i8>
- // CHECK: llvm.intr.vacopy %[[CAST0]] to %[[CAST1]]
- llvm.intr.vacopy %3 to %5 : !llvm.ptr<i8>, !llvm.ptr<i8>
- // CHECK: llvm.intr.vaend %[[CAST1]]
- // CHECK: llvm.intr.vaend %[[CAST0]]
- llvm.intr.vaend %5 : !llvm.ptr<i8>
- llvm.intr.vaend %3 : !llvm.ptr<i8>
- // CHECK: llvm.return
- llvm.return
-}
diff --git a/mlir/test/Dialect/LLVMIR/sroa.mlir b/mlir/test/Dialect/LLVMIR/sroa.mlir
index d7bf7942686ff9c..09ca94e91d2a12b 100644
--- a/mlir/test/Dialect/LLVMIR/sroa.mlir
+++ b/mlir/test/Dialect/LLVMIR/sroa.mlir
@@ -195,17 +195,3 @@ llvm.func @no_dynamic_indexing(%arg: i32) -> i32 {
// CHECK: llvm.return %[[RES]] : i32
llvm.return %3 : i32
}
-
-// -----
-
-// CHECK-LABEL: llvm.func @no_typed_pointers
-llvm.func @no_typed_pointers() -> i32 {
- // CHECK: %[[SIZE:.*]] = llvm.mlir.constant(1 : i32)
- %0 = llvm.mlir.constant(1 : i32) : i32
- // CHECK: %[[ALLOCA:.*]] = llvm.alloca %[[SIZE]] x !llvm.array<10 x i32> {alignment = 8 : i64} : (i32) -> !llvm.ptr<array<10 x i32>>
- %1 = llvm.alloca %0 x !llvm.array<10 x i32> {alignment = 8 : i64} : (i32) -> !llvm.ptr<array<10 x i32>>
- // CHECK-NOT: = llvm.alloca
- %2 = llvm.getelementptr %1[0, 1] : (!llvm.ptr<array<10 x i32>>) -> !llvm.ptr<i32>
- %3 = llvm.load %2 : !llvm.ptr<i32>
- llvm.return %3 : i32
-}
diff --git a/mlir/test/Dialect/LLVMIR/types-invalid-typed-pointers.mlir b/mlir/test/Dialect/LLVMIR/types-invalid-typed-pointers.mlir
deleted file mode 100644
index 475fadede8fbf30..000000000000000
--- a/mlir/test/Dialect/LLVMIR/types-invalid-typed-pointers.mlir
+++ /dev/null
@@ -1,42 +0,0 @@
-// 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 different 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-typed-pointers.mlir b/mlir/test/Dialect/LLVMIR/types-typed-pointers.mlir
deleted file mode 100644
index 2d63f379c2ee735..000000000000000
--- a/mlir/test/Dialect/LLVMIR/types-typed-pointers.mlir
+++ /dev/null
@@ -1,118 +0,0 @@
-// 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 c9bce337a3b8a90..2dd292408fa60dd 100644
--- a/mlir/test/Dialect/LLVMIR/types.mlir
+++ b/mlir/test/Dialect/LLVMIR/types.mlir
@@ -63,8 +63,8 @@ func.func @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>
+ // CHECK: !llvm.ptr<9>
+ "some.op"() : () -> !llvm.ptr<9>
return
}
diff --git a/mlir/test/Dialect/OpenACC/invalid.mlir b/mlir/test/Dialect/OpenACC/invalid.mlir
index b5241a8e4dc47fa..b9ac68d0592c872 100644
--- a/mlir/test/Dialect/OpenACC/invalid.mlir
+++ b/mlir/test/Dialect/OpenACC/invalid.mlir
@@ -262,31 +262,29 @@ acc.kernels dataOperands(%value : memref<10xf32>) {
// -----
// expected-error at +1 {{expects non-empty init region}}
-acc.private.recipe @privatization_i32 : !llvm.ptr<i32> init {
+acc.private.recipe @privatization_i32 : !llvm.ptr init {
}
// -----
// expected-error at +1 {{expects init region first argument of the privatization type}}
-acc.private.recipe @privatization_i32 : !llvm.ptr<i32> init {
-^bb0(%arg0 : !llvm.ptr<f32>):
+acc.private.recipe @privatization_i32 : !llvm.ptr init {
+^bb0(%arg0 : i32):
%c1 = arith.constant 1 : i32
- %c0 = arith.constant 0 : i32
- %0 = llvm.alloca %c1 x i32 : (i32) -> !llvm.ptr<i32>
- llvm.store %c0, %0 : !llvm.ptr<i32>
- acc.yield %0 : !llvm.ptr<i32>
+ %0 = llvm.alloca %c1 x i32 : (i32) -> !llvm.ptr
+ acc.yield %0 : !llvm.ptr
}
// -----
// expected-error at +1 {{expects destroy region first argument of the privatization type}}
-acc.private.recipe @privatization_i32 : !llvm.ptr<i32> init {
-^bb0(%arg0 : !llvm.ptr<i32>):
+acc.private.recipe @privatization_i32 : !llvm.ptr init {
+^bb0(%arg0 : !llvm.ptr):
%c1 = arith.constant 1 : i32
%c0 = arith.constant 0 : i32
- %0 = llvm.alloca %c1 x i32 : (i32) -> !llvm.ptr<i32>
- llvm.store %c0, %0 : !llvm.ptr<i32>
- acc.yield %0 : !llvm.ptr<i32>
+ %0 = llvm.alloca %c1 x i32 : (i32) -> !llvm.ptr
+ llvm.store %c0, %0 : i32, !llvm.ptr
+ acc.yield %0 : !llvm.ptr
} destroy {
^bb0(%arg0 : f32):
"test.openacc_dummy_op"(%arg0) : (f32) -> ()
@@ -295,44 +293,42 @@ acc.private.recipe @privatization_i32 : !llvm.ptr<i32> init {
// -----
// expected-error at +1 {{expects non-empty init region}}
-acc.firstprivate.recipe @privatization_i32 : !llvm.ptr<i32> init {
+acc.firstprivate.recipe @privatization_i32 : !llvm.ptr init {
} copy {}
// -----
// expected-error at +1 {{expects init region first argument of the privatization type}}
-acc.firstprivate.recipe @privatization_i32 : !llvm.ptr<i32> init {
-^bb0(%arg0 : !llvm.ptr<f32>):
+acc.firstprivate.recipe @privatization_i32 : !llvm.ptr init {
+^bb0(%arg0 : i32):
%c1 = arith.constant 1 : i32
- %c0 = arith.constant 0 : i32
- %0 = llvm.alloca %c1 x i32 : (i32) -> !llvm.ptr<i32>
- llvm.store %c0, %0 : !llvm.ptr<i32>
- acc.yield %0 : !llvm.ptr<i32>
+ %0 = llvm.alloca %c1 x i32 : (i32) -> !llvm.ptr
+ acc.yield %0 : !llvm.ptr
} copy {}
// -----
// expected-error at +1 {{expects non-empty copy region}}
-acc.firstprivate.recipe @privatization_i32 : !llvm.ptr<i32> init {
-^bb0(%arg0 : !llvm.ptr<i32>):
+acc.firstprivate.recipe @privatization_i32 : !llvm.ptr init {
+^bb0(%arg0 : !llvm.ptr):
%c1 = arith.constant 1 : i32
%c0 = arith.constant 0 : i32
- %0 = llvm.alloca %c1 x i32 : (i32) -> !llvm.ptr<i32>
- llvm.store %c0, %0 : !llvm.ptr<i32>
- acc.yield %0 : !llvm.ptr<i32>
+ %0 = llvm.alloca %c1 x i32 : (i32) -> !llvm.ptr
+ llvm.store %c0, %0 : i32, !llvm.ptr
+ acc.yield %0 : !llvm.ptr
} copy {
}
// -----
// expected-error at +1 {{expects copy region with two arguments of the privatization type}}
-acc.firstprivate.recipe @privatization_i32 : !llvm.ptr<i32> init {
-^bb0(%arg0 : !llvm.ptr<i32>):
+acc.firstprivate.recipe @privatization_i32 : !llvm.ptr init {
+^bb0(%arg0 : !llvm.ptr):
%c1 = arith.constant 1 : i32
%c0 = arith.constant 0 : i32
- %0 = llvm.alloca %c1 x i32 : (i32) -> !llvm.ptr<i32>
- llvm.store %c0, %0 : !llvm.ptr<i32>
- acc.yield %0 : !llvm.ptr<i32>
+ %0 = llvm.alloca %c1 x i32 : (i32) -> !llvm.ptr
+ llvm.store %c0, %0 : i32, !llvm.ptr
+ acc.yield %0 : !llvm.ptr
} copy {
^bb0(%arg0 : f32):
"test.openacc_dummy_op"(%arg0) : (f32) -> ()
@@ -341,13 +337,13 @@ acc.firstprivate.recipe @privatization_i32 : !llvm.ptr<i32> init {
// -----
// expected-error at +1 {{expects copy region with two arguments of the privatization type}}
-acc.firstprivate.recipe @privatization_i32 : !llvm.ptr<i32> init {
-^bb0(%arg0 : !llvm.ptr<i32>):
+acc.firstprivate.recipe @privatization_i32 : !llvm.ptr init {
+^bb0(%arg0 : !llvm.ptr):
%c1 = arith.constant 1 : i32
%c0 = arith.constant 0 : i32
- %0 = llvm.alloca %c1 x i32 : (i32) -> !llvm.ptr<i32>
- llvm.store %c0, %0 : !llvm.ptr<i32>
- acc.yield %0 : !llvm.ptr<i32>
+ %0 = llvm.alloca %c1 x i32 : (i32) -> !llvm.ptr
+ llvm.store %c0, %0 : i32, !llvm.ptr
+ acc.yield %0 : !llvm.ptr
} copy {
^bb0(%arg0 : f32, %arg1 : i32):
"test.openacc_dummy_op"(%arg0) : (f32) -> ()
@@ -361,8 +357,8 @@ acc.firstprivate.recipe @privatization_i32 : i32 init {
%0 = arith.constant 1 : i32
acc.yield %0 : i32
} copy {
-^bb0(%arg0 : i32, %arg1 : !llvm.ptr<i32>):
- llvm.store %arg0, %arg1 : !llvm.ptr<i32>
+^bb0(%arg0 : i32, %arg1 : !llvm.ptr):
+ llvm.store %arg0, %arg1 : i32, !llvm.ptr
acc.yield
} destroy {
^bb0(%arg0 : f32):
@@ -448,9 +444,9 @@ acc.loop gang(static=%i64Value: i64, ) {
// -----
-func.func @fct1(%0 : !llvm.ptr<i32>) -> () {
+func.func @fct1(%0 : !llvm.ptr) -> () {
// expected-error at +1 {{expected symbol reference @privatization_i32 to point to a private declaration}}
- acc.serial private(@privatization_i32 -> %0 : !llvm.ptr<i32>) {
+ acc.serial private(@privatization_i32 -> %0 : !llvm.ptr) {
}
return
}
diff --git a/mlir/test/Dialect/OpenACC/ops.mlir b/mlir/test/Dialect/OpenACC/ops.mlir
index cf7a838f55ef855..c06e5a596923f94 100644
--- a/mlir/test/Dialect/OpenACC/ops.mlir
+++ b/mlir/test/Dialect/OpenACC/ops.mlir
@@ -1101,159 +1101,159 @@ acc.shutdown device_num(%idxValue : index) if(%ifCond)
// -----
-func.func @testexitdataop(%a: !llvm.ptr<f32>) -> () {
+func.func @testexitdataop(%a: !llvm.ptr) -> () {
%ifCond = arith.constant true
%i64Value = arith.constant 1 : i64
%i32Value = arith.constant 1 : i32
%idxValue = arith.constant 1 : index
- %0 = acc.getdeviceptr varPtr(%a : !llvm.ptr<f32>) -> !llvm.ptr<f32>
- acc.exit_data dataOperands(%0 : !llvm.ptr<f32>)
- acc.copyout accPtr(%0 : !llvm.ptr<f32>) to varPtr(%a : !llvm.ptr<f32>)
+ %0 = acc.getdeviceptr varPtr(%a : !llvm.ptr) -> !llvm.ptr
+ acc.exit_data dataOperands(%0 : !llvm.ptr)
+ acc.copyout accPtr(%0 : !llvm.ptr) to varPtr(%a : !llvm.ptr)
- %1 = acc.getdeviceptr varPtr(%a : !llvm.ptr<f32>) -> !llvm.ptr<f32>
- acc.exit_data dataOperands(%1 : !llvm.ptr<f32>)
- acc.delete accPtr(%1 : !llvm.ptr<f32>)
+ %1 = acc.getdeviceptr varPtr(%a : !llvm.ptr) -> !llvm.ptr
+ acc.exit_data dataOperands(%1 : !llvm.ptr)
+ acc.delete accPtr(%1 : !llvm.ptr)
- %2 = acc.getdeviceptr varPtr(%a : !llvm.ptr<f32>) -> !llvm.ptr<f32>
- acc.exit_data dataOperands(%2 : !llvm.ptr<f32>) attributes {async,finalize}
- acc.delete accPtr(%2 : !llvm.ptr<f32>)
+ %2 = acc.getdeviceptr varPtr(%a : !llvm.ptr) -> !llvm.ptr
+ acc.exit_data dataOperands(%2 : !llvm.ptr) attributes {async,finalize}
+ acc.delete accPtr(%2 : !llvm.ptr)
- %3 = acc.getdeviceptr varPtr(%a : !llvm.ptr<f32>) -> !llvm.ptr<f32>
- acc.exit_data dataOperands(%3 : !llvm.ptr<f32>)
- acc.detach accPtr(%3 : !llvm.ptr<f32>)
+ %3 = acc.getdeviceptr varPtr(%a : !llvm.ptr) -> !llvm.ptr
+ acc.exit_data dataOperands(%3 : !llvm.ptr)
+ acc.detach accPtr(%3 : !llvm.ptr)
- %4 = acc.getdeviceptr varPtr(%a : !llvm.ptr<f32>) -> !llvm.ptr<f32>
- acc.exit_data dataOperands(%4 : !llvm.ptr<f32>) attributes {async}
- acc.copyout accPtr(%4 : !llvm.ptr<f32>) to varPtr(%a : !llvm.ptr<f32>)
+ %4 = acc.getdeviceptr varPtr(%a : !llvm.ptr) -> !llvm.ptr
+ acc.exit_data dataOperands(%4 : !llvm.ptr) attributes {async}
+ acc.copyout accPtr(%4 : !llvm.ptr) to varPtr(%a : !llvm.ptr)
- %5 = acc.getdeviceptr varPtr(%a : !llvm.ptr<f32>) -> !llvm.ptr<f32>
- acc.exit_data dataOperands(%5 : !llvm.ptr<f32>) attributes {wait}
- acc.delete accPtr(%5 : !llvm.ptr<f32>)
+ %5 = acc.getdeviceptr varPtr(%a : !llvm.ptr) -> !llvm.ptr
+ acc.exit_data dataOperands(%5 : !llvm.ptr) attributes {wait}
+ acc.delete accPtr(%5 : !llvm.ptr)
- %6 = acc.getdeviceptr varPtr(%a : !llvm.ptr<f32>) -> !llvm.ptr<f32>
- acc.exit_data async(%i64Value : i64) dataOperands(%6 : !llvm.ptr<f32>)
- acc.copyout accPtr(%6 : !llvm.ptr<f32>) to varPtr(%a : !llvm.ptr<f32>)
+ %6 = acc.getdeviceptr varPtr(%a : !llvm.ptr) -> !llvm.ptr
+ acc.exit_data async(%i64Value : i64) dataOperands(%6 : !llvm.ptr)
+ acc.copyout accPtr(%6 : !llvm.ptr) to varPtr(%a : !llvm.ptr)
- %7 = acc.getdeviceptr varPtr(%a : !llvm.ptr<f32>) -> !llvm.ptr<f32>
- acc.exit_data dataOperands(%7 : !llvm.ptr<f32>) async(%i64Value : i64)
- acc.copyout accPtr(%7 : !llvm.ptr<f32>) to varPtr(%a : !llvm.ptr<f32>)
+ %7 = acc.getdeviceptr varPtr(%a : !llvm.ptr) -> !llvm.ptr
+ acc.exit_data dataOperands(%7 : !llvm.ptr) async(%i64Value : i64)
+ acc.copyout accPtr(%7 : !llvm.ptr) to varPtr(%a : !llvm.ptr)
- %8 = acc.getdeviceptr varPtr(%a : !llvm.ptr<f32>) -> !llvm.ptr<f32>
- acc.exit_data if(%ifCond) dataOperands(%8 : !llvm.ptr<f32>)
- acc.copyout accPtr(%8 : !llvm.ptr<f32>) to varPtr(%a : !llvm.ptr<f32>)
+ %8 = acc.getdeviceptr varPtr(%a : !llvm.ptr) -> !llvm.ptr
+ acc.exit_data if(%ifCond) dataOperands(%8 : !llvm.ptr)
+ acc.copyout accPtr(%8 : !llvm.ptr) to varPtr(%a : !llvm.ptr)
- %9 = acc.getdeviceptr varPtr(%a : !llvm.ptr<f32>) -> !llvm.ptr<f32>
- acc.exit_data wait_devnum(%i64Value: i64) wait(%i32Value, %idxValue : i32, index) dataOperands(%9 : !llvm.ptr<f32>)
- acc.copyout accPtr(%9 : !llvm.ptr<f32>) to varPtr(%a : !llvm.ptr<f32>)
+ %9 = acc.getdeviceptr varPtr(%a : !llvm.ptr) -> !llvm.ptr
+ acc.exit_data wait_devnum(%i64Value: i64) wait(%i32Value, %idxValue : i32, index) dataOperands(%9 : !llvm.ptr)
+ acc.copyout accPtr(%9 : !llvm.ptr) to varPtr(%a : !llvm.ptr)
return
}
-// CHECK: func @testexitdataop(%[[ARGA:.*]]: !llvm.ptr<f32>) {
+// CHECK: func @testexitdataop(%[[ARGA:.*]]: !llvm.ptr) {
// CHECK: %[[IFCOND:.*]] = arith.constant true
// CHECK: %[[I64VALUE:.*]] = arith.constant 1 : i64
// CHECK: %[[I32VALUE:.*]] = arith.constant 1 : i32
// CHECK: %[[IDXVALUE:.*]] = arith.constant 1 : index
-// CHECK: %[[DEVPTR:.*]] = acc.getdeviceptr varPtr(%[[ARGA]] : !llvm.ptr<f32>) -> !llvm.ptr<f32>
-// CHECK: acc.exit_data dataOperands(%[[DEVPTR]] : !llvm.ptr<f32>)
-// CHECK: acc.copyout accPtr(%[[DEVPTR]] : !llvm.ptr<f32>) to varPtr(%[[ARGA]] : !llvm.ptr<f32>)
+// CHECK: %[[DEVPTR:.*]] = acc.getdeviceptr varPtr(%[[ARGA]] : !llvm.ptr) -> !llvm.ptr
+// CHECK: acc.exit_data dataOperands(%[[DEVPTR]] : !llvm.ptr)
+// CHECK: acc.copyout accPtr(%[[DEVPTR]] : !llvm.ptr) to varPtr(%[[ARGA]] : !llvm.ptr)
-// CHECK: %[[DEVPTR:.*]] = acc.getdeviceptr varPtr(%[[ARGA]] : !llvm.ptr<f32>) -> !llvm.ptr<f32>
-// CHECK: acc.exit_data dataOperands(%[[DEVPTR]] : !llvm.ptr<f32>)
-// CHECK: acc.delete accPtr(%[[DEVPTR]] : !llvm.ptr<f32>)
+// CHECK: %[[DEVPTR:.*]] = acc.getdeviceptr varPtr(%[[ARGA]] : !llvm.ptr) -> !llvm.ptr
+// CHECK: acc.exit_data dataOperands(%[[DEVPTR]] : !llvm.ptr)
+// CHECK: acc.delete accPtr(%[[DEVPTR]] : !llvm.ptr)
-// CHECK: %[[DEVPTR:.*]] = acc.getdeviceptr varPtr(%[[ARGA]] : !llvm.ptr<f32>) -> !llvm.ptr<f32>
-// CHECK: acc.exit_data dataOperands(%[[DEVPTR]] : !llvm.ptr<f32>) attributes {async, finalize}
-// CHECK: acc.delete accPtr(%[[DEVPTR]] : !llvm.ptr<f32>)
+// CHECK: %[[DEVPTR:.*]] = acc.getdeviceptr varPtr(%[[ARGA]] : !llvm.ptr) -> !llvm.ptr
+// CHECK: acc.exit_data dataOperands(%[[DEVPTR]] : !llvm.ptr) attributes {async, finalize}
+// CHECK: acc.delete accPtr(%[[DEVPTR]] : !llvm.ptr)
-// CHECK: %[[DEVPTR:.*]] = acc.getdeviceptr varPtr(%[[ARGA]] : !llvm.ptr<f32>) -> !llvm.ptr<f32>
-// CHECK: acc.exit_data dataOperands(%[[DEVPTR]] : !llvm.ptr<f32>)
-// CHECK: acc.detach accPtr(%[[DEVPTR]] : !llvm.ptr<f32>)
+// CHECK: %[[DEVPTR:.*]] = acc.getdeviceptr varPtr(%[[ARGA]] : !llvm.ptr) -> !llvm.ptr
+// CHECK: acc.exit_data dataOperands(%[[DEVPTR]] : !llvm.ptr)
+// CHECK: acc.detach accPtr(%[[DEVPTR]] : !llvm.ptr)
-// CHECK: %[[DEVPTR:.*]] = acc.getdeviceptr varPtr(%[[ARGA]] : !llvm.ptr<f32>) -> !llvm.ptr<f32>
-// CHECK: acc.exit_data dataOperands(%[[DEVPTR]] : !llvm.ptr<f32>) attributes {async}
-// CHECK: acc.copyout accPtr(%[[DEVPTR]] : !llvm.ptr<f32>) to varPtr(%[[ARGA]] : !llvm.ptr<f32>)
+// CHECK: %[[DEVPTR:.*]] = acc.getdeviceptr varPtr(%[[ARGA]] : !llvm.ptr) -> !llvm.ptr
+// CHECK: acc.exit_data dataOperands(%[[DEVPTR]] : !llvm.ptr) attributes {async}
+// CHECK: acc.copyout accPtr(%[[DEVPTR]] : !llvm.ptr) to varPtr(%[[ARGA]] : !llvm.ptr)
-// CHECK: %[[DEVPTR:.*]] = acc.getdeviceptr varPtr(%[[ARGA]] : !llvm.ptr<f32>) -> !llvm.ptr<f32>
-// CHECK: acc.exit_data dataOperands(%[[DEVPTR]] : !llvm.ptr<f32>) attributes {wait}
-// CHECK: acc.delete accPtr(%[[DEVPTR]] : !llvm.ptr<f32>)
+// CHECK: %[[DEVPTR:.*]] = acc.getdeviceptr varPtr(%[[ARGA]] : !llvm.ptr) -> !llvm.ptr
+// CHECK: acc.exit_data dataOperands(%[[DEVPTR]] : !llvm.ptr) attributes {wait}
+// CHECK: acc.delete accPtr(%[[DEVPTR]] : !llvm.ptr)
-// CHECK: %[[DEVPTR:.*]] = acc.getdeviceptr varPtr(%[[ARGA]] : !llvm.ptr<f32>) -> !llvm.ptr<f32>
-// CHECK: acc.exit_data async(%[[I64VALUE]] : i64) dataOperands(%[[DEVPTR]] : !llvm.ptr<f32>)
-// CHECK: acc.copyout accPtr(%[[DEVPTR]] : !llvm.ptr<f32>) to varPtr(%[[ARGA]] : !llvm.ptr<f32>)
+// CHECK: %[[DEVPTR:.*]] = acc.getdeviceptr varPtr(%[[ARGA]] : !llvm.ptr) -> !llvm.ptr
+// CHECK: acc.exit_data async(%[[I64VALUE]] : i64) dataOperands(%[[DEVPTR]] : !llvm.ptr)
+// CHECK: acc.copyout accPtr(%[[DEVPTR]] : !llvm.ptr) to varPtr(%[[ARGA]] : !llvm.ptr)
-// CHECK: %[[DEVPTR:.*]] = acc.getdeviceptr varPtr(%[[ARGA]] : !llvm.ptr<f32>) -> !llvm.ptr<f32>
-// CHECK: acc.exit_data async(%[[I64VALUE]] : i64) dataOperands(%[[DEVPTR]] : !llvm.ptr<f32>)
-// CHECK: acc.copyout accPtr(%[[DEVPTR]] : !llvm.ptr<f32>) to varPtr(%[[ARGA]] : !llvm.ptr<f32>)
+// CHECK: %[[DEVPTR:.*]] = acc.getdeviceptr varPtr(%[[ARGA]] : !llvm.ptr) -> !llvm.ptr
+// CHECK: acc.exit_data async(%[[I64VALUE]] : i64) dataOperands(%[[DEVPTR]] : !llvm.ptr)
+// CHECK: acc.copyout accPtr(%[[DEVPTR]] : !llvm.ptr) to varPtr(%[[ARGA]] : !llvm.ptr)
-// CHECK: %[[DEVPTR:.*]] = acc.getdeviceptr varPtr(%[[ARGA]] : !llvm.ptr<f32>) -> !llvm.ptr<f32>
-// CHECK: acc.exit_data if(%[[IFCOND]]) dataOperands(%[[DEVPTR]] : !llvm.ptr<f32>)
-// CHECK: acc.copyout accPtr(%[[DEVPTR]] : !llvm.ptr<f32>) to varPtr(%[[ARGA]] : !llvm.ptr<f32>)
+// CHECK: %[[DEVPTR:.*]] = acc.getdeviceptr varPtr(%[[ARGA]] : !llvm.ptr) -> !llvm.ptr
+// CHECK: acc.exit_data if(%[[IFCOND]]) dataOperands(%[[DEVPTR]] : !llvm.ptr)
+// CHECK: acc.copyout accPtr(%[[DEVPTR]] : !llvm.ptr) to varPtr(%[[ARGA]] : !llvm.ptr)
-// CHECK: %[[DEVPTR:.*]] = acc.getdeviceptr varPtr(%[[ARGA]] : !llvm.ptr<f32>) -> !llvm.ptr<f32>
-// CHECK: acc.exit_data wait_devnum(%[[I64VALUE]] : i64) wait(%[[I32VALUE]], %[[IDXVALUE]] : i32, index) dataOperands(%[[DEVPTR]] : !llvm.ptr<f32>)
-// CHECK: acc.copyout accPtr(%[[DEVPTR]] : !llvm.ptr<f32>) to varPtr(%[[ARGA]] : !llvm.ptr<f32>)
+// CHECK: %[[DEVPTR:.*]] = acc.getdeviceptr varPtr(%[[ARGA]] : !llvm.ptr) -> !llvm.ptr
+// CHECK: acc.exit_data wait_devnum(%[[I64VALUE]] : i64) wait(%[[I32VALUE]], %[[IDXVALUE]] : i32, index) dataOperands(%[[DEVPTR]] : !llvm.ptr)
+// CHECK: acc.copyout accPtr(%[[DEVPTR]] : !llvm.ptr) to varPtr(%[[ARGA]] : !llvm.ptr)
// -----
-func.func @testenterdataop(%a: !llvm.ptr<f32>, %b: !llvm.ptr<f32>, %c: !llvm.ptr<f32>) -> () {
+func.func @testenterdataop(%a: !llvm.ptr, %b: !llvm.ptr, %c: !llvm.ptr) -> () {
%ifCond = arith.constant true
%i64Value = arith.constant 1 : i64
%i32Value = arith.constant 1 : i32
%idxValue = arith.constant 1 : index
- %0 = acc.copyin varPtr(%a : !llvm.ptr<f32>) -> !llvm.ptr<f32>
- acc.enter_data dataOperands(%0 : !llvm.ptr<f32>)
- %1 = acc.create varPtr(%a : !llvm.ptr<f32>) -> !llvm.ptr<f32>
- %2 = acc.create varPtr(%b : !llvm.ptr<f32>) -> !llvm.ptr<f32> {dataClause = #acc<data_clause acc_create_zero>}
- %3 = acc.create varPtr(%c : !llvm.ptr<f32>) -> !llvm.ptr<f32> {dataClause = #acc<data_clause acc_create_zero>}
- acc.enter_data dataOperands(%1, %2, %3 : !llvm.ptr<f32>, !llvm.ptr<f32>, !llvm.ptr<f32>)
- %4 = acc.attach varPtr(%a : !llvm.ptr<f32>) -> !llvm.ptr<f32>
- acc.enter_data dataOperands(%4 : !llvm.ptr<f32>)
- %5 = acc.copyin varPtr(%a : !llvm.ptr<f32>) -> !llvm.ptr<f32>
- acc.enter_data dataOperands(%5 : !llvm.ptr<f32>) attributes {async}
- %6 = acc.create varPtr(%a : !llvm.ptr<f32>) -> !llvm.ptr<f32>
- acc.enter_data dataOperands(%6 : !llvm.ptr<f32>) attributes {wait}
- %7 = acc.copyin varPtr(%a : !llvm.ptr<f32>) -> !llvm.ptr<f32>
- acc.enter_data async(%i64Value : i64) dataOperands(%7 : !llvm.ptr<f32>)
- %8 = acc.copyin varPtr(%a : !llvm.ptr<f32>) -> !llvm.ptr<f32>
- acc.enter_data dataOperands(%8 : !llvm.ptr<f32>) async(%i64Value : i64)
- %9 = acc.copyin varPtr(%a : !llvm.ptr<f32>) -> !llvm.ptr<f32>
- acc.enter_data if(%ifCond) dataOperands(%9 : !llvm.ptr<f32>)
- %10 = acc.copyin varPtr(%a : !llvm.ptr<f32>) -> !llvm.ptr<f32>
- acc.enter_data wait_devnum(%i64Value: i64) wait(%i32Value, %idxValue : i32, index) dataOperands(%10 : !llvm.ptr<f32>)
+ %0 = acc.copyin varPtr(%a : !llvm.ptr) -> !llvm.ptr
+ acc.enter_data dataOperands(%0 : !llvm.ptr)
+ %1 = acc.create varPtr(%a : !llvm.ptr) -> !llvm.ptr
+ %2 = acc.create varPtr(%b : !llvm.ptr) -> !llvm.ptr {dataClause = #acc<data_clause acc_create_zero>}
+ %3 = acc.create varPtr(%c : !llvm.ptr) -> !llvm.ptr {dataClause = #acc<data_clause acc_create_zero>}
+ acc.enter_data dataOperands(%1, %2, %3 : !llvm.ptr, !llvm.ptr, !llvm.ptr)
+ %4 = acc.attach varPtr(%a : !llvm.ptr) -> !llvm.ptr
+ acc.enter_data dataOperands(%4 : !llvm.ptr)
+ %5 = acc.copyin varPtr(%a : !llvm.ptr) -> !llvm.ptr
+ acc.enter_data dataOperands(%5 : !llvm.ptr) attributes {async}
+ %6 = acc.create varPtr(%a : !llvm.ptr) -> !llvm.ptr
+ acc.enter_data dataOperands(%6 : !llvm.ptr) attributes {wait}
+ %7 = acc.copyin varPtr(%a : !llvm.ptr) -> !llvm.ptr
+ acc.enter_data async(%i64Value : i64) dataOperands(%7 : !llvm.ptr)
+ %8 = acc.copyin varPtr(%a : !llvm.ptr) -> !llvm.ptr
+ acc.enter_data dataOperands(%8 : !llvm.ptr) async(%i64Value : i64)
+ %9 = acc.copyin varPtr(%a : !llvm.ptr) -> !llvm.ptr
+ acc.enter_data if(%ifCond) dataOperands(%9 : !llvm.ptr)
+ %10 = acc.copyin varPtr(%a : !llvm.ptr) -> !llvm.ptr
+ acc.enter_data wait_devnum(%i64Value: i64) wait(%i32Value, %idxValue : i32, index) dataOperands(%10 : !llvm.ptr)
return
}
-// CHECK: func @testenterdataop(%[[ARGA:.*]]: !llvm.ptr<f32>, %[[ARGB:.*]]: !llvm.ptr<f32>, %[[ARGC:.*]]: !llvm.ptr<f32>) {
+// CHECK: func @testenterdataop(%[[ARGA:.*]]: !llvm.ptr, %[[ARGB:.*]]: !llvm.ptr, %[[ARGC:.*]]: !llvm.ptr) {
// CHECK: [[IFCOND:%.*]] = arith.constant true
// CHECK: [[I64VALUE:%.*]] = arith.constant 1 : i64
// CHECK: [[I32VALUE:%.*]] = arith.constant 1 : i32
// CHECK: [[IDXVALUE:%.*]] = arith.constant 1 : index
-// CHECK: %[[COPYIN:.*]] = acc.copyin varPtr(%[[ARGA]] : !llvm.ptr<f32>) -> !llvm.ptr<f32>
-// CHECK: acc.enter_data dataOperands(%[[COPYIN]] : !llvm.ptr<f32>)
-// CHECK: %[[CREATE_A:.*]] = acc.create varPtr(%[[ARGA]] : !llvm.ptr<f32>) -> !llvm.ptr<f32>
-// CHECK: %[[CREATE_B:.*]] = acc.create varPtr(%[[ARGB]] : !llvm.ptr<f32>) -> !llvm.ptr<f32> {dataClause = #acc<data_clause acc_create_zero>}
-// CHECK: %[[CREATE_C:.*]] = acc.create varPtr(%[[ARGC]] : !llvm.ptr<f32>) -> !llvm.ptr<f32> {dataClause = #acc<data_clause acc_create_zero>}
-// CHECK: acc.enter_data dataOperands(%[[CREATE_A]], %[[CREATE_B]], %[[CREATE_C]] : !llvm.ptr<f32>, !llvm.ptr<f32>, !llvm.ptr<f32>)
-// CHECK: %[[ATTACH:.*]] = acc.attach varPtr(%[[ARGA]] : !llvm.ptr<f32>) -> !llvm.ptr<f32>
-// CHECK: acc.enter_data dataOperands(%[[ATTACH]] : !llvm.ptr<f32>)
-// CHECK: %[[COPYIN:.*]] = acc.copyin varPtr(%[[ARGA]] : !llvm.ptr<f32>) -> !llvm.ptr<f32>
-// CHECK: acc.enter_data dataOperands(%[[COPYIN]] : !llvm.ptr<f32>) attributes {async}
-// CHECK: %[[CREATE:.*]] = acc.create varPtr(%[[ARGA]] : !llvm.ptr<f32>) -> !llvm.ptr<f32>
-// CHECK: acc.enter_data dataOperands(%[[CREATE]] : !llvm.ptr<f32>) attributes {wait}
-// CHECK: %[[COPYIN:.*]] = acc.copyin varPtr(%[[ARGA]] : !llvm.ptr<f32>) -> !llvm.ptr<f32>
-// CHECK: acc.enter_data async([[I64VALUE]] : i64) dataOperands(%[[COPYIN]] : !llvm.ptr<f32>)
-// CHECK: %[[COPYIN:.*]] = acc.copyin varPtr(%[[ARGA]] : !llvm.ptr<f32>) -> !llvm.ptr<f32>
-// CHECK: acc.enter_data async([[I64VALUE]] : i64) dataOperands(%[[COPYIN]] : !llvm.ptr<f32>)
-// CHECK: %[[COPYIN:.*]] = acc.copyin varPtr(%[[ARGA]] : !llvm.ptr<f32>) -> !llvm.ptr<f32>
-// CHECK: acc.enter_data if([[IFCOND]]) dataOperands(%[[COPYIN]] : !llvm.ptr<f32>)
-// CHECK: %[[COPYIN:.*]] = acc.copyin varPtr(%[[ARGA]] : !llvm.ptr<f32>) -> !llvm.ptr<f32>
-// CHECK: acc.enter_data wait_devnum([[I64VALUE]] : i64) wait([[I32VALUE]], [[IDXVALUE]] : i32, index) dataOperands(%[[COPYIN]] : !llvm.ptr<f32>)
+// CHECK: %[[COPYIN:.*]] = acc.copyin varPtr(%[[ARGA]] : !llvm.ptr) -> !llvm.ptr
+// CHECK: acc.enter_data dataOperands(%[[COPYIN]] : !llvm.ptr)
+// CHECK: %[[CREATE_A:.*]] = acc.create varPtr(%[[ARGA]] : !llvm.ptr) -> !llvm.ptr
+// CHECK: %[[CREATE_B:.*]] = acc.create varPtr(%[[ARGB]] : !llvm.ptr) -> !llvm.ptr {dataClause = #acc<data_clause acc_create_zero>}
+// CHECK: %[[CREATE_C:.*]] = acc.create varPtr(%[[ARGC]] : !llvm.ptr) -> !llvm.ptr {dataClause = #acc<data_clause acc_create_zero>}
+// CHECK: acc.enter_data dataOperands(%[[CREATE_A]], %[[CREATE_B]], %[[CREATE_C]] : !llvm.ptr, !llvm.ptr, !llvm.ptr)
+// CHECK: %[[ATTACH:.*]] = acc.attach varPtr(%[[ARGA]] : !llvm.ptr) -> !llvm.ptr
+// CHECK: acc.enter_data dataOperands(%[[ATTACH]] : !llvm.ptr)
+// CHECK: %[[COPYIN:.*]] = acc.copyin varPtr(%[[ARGA]] : !llvm.ptr) -> !llvm.ptr
+// CHECK: acc.enter_data dataOperands(%[[COPYIN]] : !llvm.ptr) attributes {async}
+// CHECK: %[[CREATE:.*]] = acc.create varPtr(%[[ARGA]] : !llvm.ptr) -> !llvm.ptr
+// CHECK: acc.enter_data dataOperands(%[[CREATE]] : !llvm.ptr) attributes {wait}
+// CHECK: %[[COPYIN:.*]] = acc.copyin varPtr(%[[ARGA]] : !llvm.ptr) -> !llvm.ptr
+// CHECK: acc.enter_data async([[I64VALUE]] : i64) dataOperands(%[[COPYIN]] : !llvm.ptr)
+// CHECK: %[[COPYIN:.*]] = acc.copyin varPtr(%[[ARGA]] : !llvm.ptr) -> !llvm.ptr
+// CHECK: acc.enter_data async([[I64VALUE]] : i64) dataOperands(%[[COPYIN]] : !llvm.ptr)
+// CHECK: %[[COPYIN:.*]] = acc.copyin varPtr(%[[ARGA]] : !llvm.ptr) -> !llvm.ptr
+// CHECK: acc.enter_data if([[IFCOND]]) dataOperands(%[[COPYIN]] : !llvm.ptr)
+// CHECK: %[[COPYIN:.*]] = acc.copyin varPtr(%[[ARGA]] : !llvm.ptr) -> !llvm.ptr
+// CHECK: acc.enter_data wait_devnum([[I64VALUE]] : i64) wait([[I32VALUE]], [[IDXVALUE]] : i32, index) dataOperands(%[[COPYIN]] : !llvm.ptr)
// -----
@@ -1419,42 +1419,42 @@ func.func @host_device_ops(%a: memref<f32>) -> () {
// -----
-func.func @host_data_ops(%a: !llvm.ptr<f32>, %ifCond: i1) -> () {
- %0 = acc.use_device varPtr(%a : !llvm.ptr<f32>) -> !llvm.ptr<f32>
- acc.host_data dataOperands(%0: !llvm.ptr<f32>) {
+func.func @host_data_ops(%a: !llvm.ptr, %ifCond: i1) -> () {
+ %0 = acc.use_device varPtr(%a : !llvm.ptr) -> !llvm.ptr
+ acc.host_data dataOperands(%0: !llvm.ptr) {
}
- acc.host_data dataOperands(%0: !llvm.ptr<f32>) {
+ acc.host_data dataOperands(%0: !llvm.ptr) {
} attributes {if_present}
- acc.host_data if(%ifCond) dataOperands(%0: !llvm.ptr<f32>) {
+ acc.host_data if(%ifCond) dataOperands(%0: !llvm.ptr) {
}
return
}
// CHECK-LABEL: func.func @host_data_ops(
-// CHECK-SAME: %[[A:.*]]: !llvm.ptr<f32>, %[[IFCOND:.*]]: i1)
-// CHECK: %[[PTR:.*]] = acc.use_device varPtr(%[[A]] : !llvm.ptr<f32>) -> !llvm.ptr<f32>
-// CHECK: acc.host_data dataOperands(%[[PTR]] : !llvm.ptr<f32>)
-// CHECK: acc.host_data dataOperands(%[[PTR]] : !llvm.ptr<f32>) {
+// CHECK-SAME: %[[A:.*]]: !llvm.ptr, %[[IFCOND:.*]]: i1)
+// CHECK: %[[PTR:.*]] = acc.use_device varPtr(%[[A]] : !llvm.ptr) -> !llvm.ptr
+// CHECK: acc.host_data dataOperands(%[[PTR]] : !llvm.ptr)
+// CHECK: acc.host_data dataOperands(%[[PTR]] : !llvm.ptr) {
// CHECK: } attributes {if_present}
-// CHECK: acc.host_data if(%[[IFCOND]]) dataOperands(%[[PTR]] : !llvm.ptr<f32>)
+// CHECK: acc.host_data if(%[[IFCOND]]) dataOperands(%[[PTR]] : !llvm.ptr)
// -----
-acc.private.recipe @privatization_i32 : !llvm.ptr<i32> init {
-^bb0(%arg0: !llvm.ptr<i32>):
+acc.private.recipe @privatization_i32 : !llvm.ptr init {
+^bb0(%arg0: !llvm.ptr):
%c1 = arith.constant 1 : i32
%c0 = arith.constant 0 : i32
- %0 = llvm.alloca %c1 x i32 : (i32) -> !llvm.ptr<i32>
- llvm.store %c0, %0 : !llvm.ptr<i32>
- acc.yield %0 : !llvm.ptr<i32>
+ %0 = llvm.alloca %c1 x i32 : (i32) -> !llvm.ptr
+ llvm.store %c0, %0 : i32, !llvm.ptr
+ acc.yield %0 : !llvm.ptr
}
-// CHECK: acc.private.recipe @privatization_i32 : !llvm.ptr<i32> init {
+// CHECK: acc.private.recipe @privatization_i32 : !llvm.ptr init {
// CHECK: %[[C1:.*]] = arith.constant 1 : i32
// CHECK: %[[C0:.*]] = arith.constant 0 : i32
-// CHECK: %[[ALLOCA:.*]] = llvm.alloca %[[C1]] x i32 : (i32) -> !llvm.ptr<i32>
-// CHECK: llvm.store %[[C0]], %[[ALLOCA]] : !llvm.ptr<i32>
-// CHECK: acc.yield %[[ALLOCA]] : !llvm.ptr<i32>
+// CHECK: %[[ALLOCA:.*]] = llvm.alloca %[[C1]] x i32 : (i32) -> !llvm.ptr
+// CHECK: llvm.store %[[C0]], %[[ALLOCA]] : i32, !llvm.ptr
+// CHECK: acc.yield %[[ALLOCA]] : !llvm.ptr
// -----
@@ -1617,29 +1617,29 @@ llvm.mlir.global external @globalvar() { acc.declare = #acc.declare<dataClause =
}
acc.global_ctor @acc_constructor {
- %0 = llvm.mlir.addressof @globalvar { acc.declare = #acc.declare<dataClause = acc_create> } : !llvm.ptr<i32>
- %1 = acc.create varPtr(%0 : !llvm.ptr<i32>) -> !llvm.ptr<i32>
- acc.declare_enter dataOperands(%1 : !llvm.ptr<i32>)
+ %0 = llvm.mlir.addressof @globalvar { acc.declare = #acc.declare<dataClause = acc_create> } : !llvm.ptr
+ %1 = acc.create varPtr(%0 : !llvm.ptr) -> !llvm.ptr
+ acc.declare_enter dataOperands(%1 : !llvm.ptr)
acc.terminator
}
acc.global_dtor @acc_destructor {
- %0 = llvm.mlir.addressof @globalvar { acc.declare = #acc.declare<dataClause = acc_create> } : !llvm.ptr<i32>
- %1 = acc.getdeviceptr varPtr(%0 : !llvm.ptr<i32>) -> !llvm.ptr<i32> {dataClause = #acc<data_clause acc_create>}
- acc.declare_exit dataOperands(%1 : !llvm.ptr<i32>)
- acc.delete accPtr(%1 : !llvm.ptr<i32>)
+ %0 = llvm.mlir.addressof @globalvar { acc.declare = #acc.declare<dataClause = acc_create> } : !llvm.ptr
+ %1 = acc.getdeviceptr varPtr(%0 : !llvm.ptr) -> !llvm.ptr {dataClause = #acc<data_clause acc_create>}
+ acc.declare_exit dataOperands(%1 : !llvm.ptr)
+ acc.delete accPtr(%1 : !llvm.ptr)
acc.terminator
}
// CHECK-LABEL: acc.global_ctor @acc_constructor
-// CHECK: %[[ADDR:.*]] = llvm.mlir.addressof @globalvar {acc.declare = #acc.declare<dataClause = acc_create>} : !llvm.ptr<i32>
-// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[ADDR]] : !llvm.ptr<i32>) -> !llvm.ptr<i32>
-// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !llvm.ptr<i32>)
+// CHECK: %[[ADDR:.*]] = llvm.mlir.addressof @globalvar {acc.declare = #acc.declare<dataClause = acc_create>} : !llvm.ptr
+// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[ADDR]] : !llvm.ptr) -> !llvm.ptr
+// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !llvm.ptr)
// CHECK: acc.global_dtor @acc_destructor
-// CHECK: %[[ADDR:.*]] = llvm.mlir.addressof @globalvar {acc.declare = #acc.declare<dataClause = acc_create>} : !llvm.ptr<i32>
-// CHECK-NEXT: %[[DELETE:.*]] = acc.getdeviceptr varPtr(%[[ADDR]] : !llvm.ptr<i32>) -> !llvm.ptr<i32> {dataClause = #acc<data_clause acc_create>}
-// CHECK-NEXT: acc.declare_exit dataOperands(%[[DELETE]] : !llvm.ptr<i32>)
-// CHECK-NEXT: acc.delete accPtr(%[[DELETE]] : !llvm.ptr<i32>)
+// CHECK: %[[ADDR:.*]] = llvm.mlir.addressof @globalvar {acc.declare = #acc.declare<dataClause = acc_create>} : !llvm.ptr
+// CHECK-NEXT: %[[DELETE:.*]] = acc.getdeviceptr varPtr(%[[ADDR]] : !llvm.ptr) -> !llvm.ptr {dataClause = #acc<data_clause acc_create>}
+// CHECK-NEXT: acc.declare_exit dataOperands(%[[DELETE]] : !llvm.ptr)
+// CHECK-NEXT: acc.delete accPtr(%[[DELETE]] : !llvm.ptr)
// -----
@@ -1691,14 +1691,14 @@ func.func @compute3(%a: memref<10x10xf32>, %b: memref<10x10xf32>, %c: memref<10x
%numWorkers = arith.constant 10 : i64
%c20 = arith.constant 20 : i32
- %alloc = llvm.alloca %c20 x i32 { acc.declare = #acc.declare<dataClause = acc_create, implicit = true> } : (i32) -> !llvm.ptr<i32>
- %createlocal = acc.create varPtr(%alloc : !llvm.ptr<i32>) -> !llvm.ptr<i32> {implicit = true}
+ %alloc = llvm.alloca %c20 x i32 { acc.declare = #acc.declare<dataClause = acc_create, implicit = true> } : (i32) -> !llvm.ptr
+ %createlocal = acc.create varPtr(%alloc : !llvm.ptr) -> !llvm.ptr {implicit = true}
%pa = acc.present varPtr(%a : memref<10x10xf32>) -> memref<10x10xf32>
%pb = acc.present varPtr(%b : memref<10x10xf32>) -> memref<10x10xf32>
%pc = acc.present varPtr(%c : memref<10xf32>) -> memref<10xf32>
%pd = acc.present varPtr(%d : memref<10xf32>) -> memref<10xf32>
- acc.declare dataOperands(%pa, %pb, %pc, %pd, %createlocal: memref<10x10xf32>, memref<10x10xf32>, memref<10xf32>, memref<10xf32>, !llvm.ptr<i32>) {
+ acc.declare dataOperands(%pa, %pb, %pc, %pd, %createlocal: memref<10x10xf32>, memref<10x10xf32>, memref<10xf32>, memref<10xf32>, !llvm.ptr) {
}
return
diff --git a/mlir/test/Dialect/OpenMP/canonicalize.mlir b/mlir/test/Dialect/OpenMP/canonicalize.mlir
index 68f5bacb1def178..8aff8f81188be50 100644
--- a/mlir/test/Dialect/OpenMP/canonicalize.mlir
+++ b/mlir/test/Dialect/OpenMP/canonicalize.mlir
@@ -129,10 +129,10 @@ func.func private @foo() -> ()
// -----
-func.func @constant_hoisting_target(%x : !llvm.ptr<i32>) {
+func.func @constant_hoisting_target(%x : !llvm.ptr) {
omp.target {
%c1 = arith.constant 10 : i32
- llvm.store %c1, %x : i32, !llvm.ptr<i32>
+ llvm.store %c1, %x : i32, !llvm.ptr
omp.terminator
}
return
diff --git a/mlir/test/Dialect/OpenMP/invalid.mlir b/mlir/test/Dialect/OpenMP/invalid.mlir
index d4106453f31ed74..6f75f2a62e64136 100644
--- a/mlir/test/Dialect/OpenMP/invalid.mlir
+++ b/mlir/test/Dialect/OpenMP/invalid.mlir
@@ -450,14 +450,14 @@ combiner {
func.func @foo(%lb : index, %ub : index, %step : index) {
%c1 = arith.constant 1 : i32
- %0 = llvm.alloca %c1 x i32 : (i32) -> !llvm.ptr<f32>
- %1 = llvm.alloca %c1 x i32 : (i32) -> !llvm.ptr<f32>
+ %0 = llvm.alloca %c1 x i32 : (i32) -> !llvm.ptr
+ %1 = llvm.alloca %c1 x i32 : (i32) -> !llvm.ptr
- omp.wsloop reduction(@add_f32 -> %0 : !llvm.ptr<f32>)
+ omp.wsloop reduction(@add_f32 -> %0 : !llvm.ptr)
for (%iv) : index = (%lb) to (%ub) step (%step) {
%2 = arith.constant 2.0 : f32
// expected-error @below {{accumulator is not used by the parent}}
- omp.reduction %2, %1 : f32, !llvm.ptr<f32>
+ omp.reduction %2, %1 : f32, !llvm.ptr
omp.yield
}
return
@@ -467,14 +467,14 @@ func.func @foo(%lb : index, %ub : index, %step : index) {
func.func @foo(%lb : index, %ub : index, %step : index) {
%c1 = arith.constant 1 : i32
- %0 = llvm.alloca %c1 x i32 : (i32) -> !llvm.ptr<f32>
- %1 = llvm.alloca %c1 x i32 : (i32) -> !llvm.ptr<f32>
+ %0 = llvm.alloca %c1 x i32 : (i32) -> !llvm.ptr
+ %1 = llvm.alloca %c1 x i32 : (i32) -> !llvm.ptr
// expected-error @below {{expected symbol reference @foo to point to a reduction declaration}}
- omp.wsloop reduction(@foo -> %0 : !llvm.ptr<f32>)
+ omp.wsloop reduction(@foo -> %0 : !llvm.ptr)
for (%iv) : index = (%lb) to (%ub) step (%step) {
%2 = arith.constant 2.0 : f32
- omp.reduction %2, %1 : f32, !llvm.ptr<f32>
+ omp.reduction %2, %1 : f32, !llvm.ptr
omp.yield
}
return
@@ -496,13 +496,13 @@ combiner {
func.func @foo(%lb : index, %ub : index, %step : index) {
%c1 = arith.constant 1 : i32
- %0 = llvm.alloca %c1 x i32 : (i32) -> !llvm.ptr<f32>
+ %0 = llvm.alloca %c1 x i32 : (i32) -> !llvm.ptr
// expected-error @below {{accumulator variable used more than once}}
- omp.wsloop reduction(@add_f32 -> %0 : !llvm.ptr<f32>, @add_f32 -> %0 : !llvm.ptr<f32>)
+ omp.wsloop reduction(@add_f32 -> %0 : !llvm.ptr, @add_f32 -> %0 : !llvm.ptr)
for (%iv) : index = (%lb) to (%ub) step (%step) {
%2 = arith.constant 2.0 : f32
- omp.reduction %2, %0 : f32, !llvm.ptr<f32>
+ omp.reduction %2, %0 : f32, !llvm.ptr
omp.yield
}
return
@@ -522,16 +522,16 @@ combiner {
omp.yield (%1 : f32)
}
atomic {
-^bb2(%arg2: !llvm.ptr<f32>, %arg3: !llvm.ptr<f32>):
- %2 = llvm.load %arg3 : !llvm.ptr<f32>
- llvm.atomicrmw fadd %arg2, %2 monotonic : !llvm.ptr<f32>, f32
+^bb2(%arg2: !llvm.ptr, %arg3: !llvm.ptr):
+ %2 = llvm.load %arg3 : !llvm.ptr -> f32
+ llvm.atomicrmw fadd %arg2, %2 monotonic : !llvm.ptr, f32
omp.yield
}
func.func @foo(%lb : index, %ub : index, %step : index, %mem : memref<1xf32>) {
%c1 = arith.constant 1 : i32
- // expected-error @below {{expected accumulator ('memref<1xf32>') to be the same type as reduction declaration ('!llvm.ptr<f32>')}}
+ // expected-error @below {{expected accumulator ('memref<1xf32>') to be the same type as reduction declaration ('!llvm.ptr')}}
omp.wsloop reduction(@add_f32 -> %mem : memref<1xf32>)
for (%iv) : index = (%lb) to (%ub) step (%step) {
%2 = arith.constant 2.0 : f32
@@ -1300,9 +1300,9 @@ func.func @omp_task_depend(%data_var: memref<i32>) {
// -----
-func.func @omp_task(%ptr: !llvm.ptr<f32>) {
+func.func @omp_task(%ptr: !llvm.ptr) {
// expected-error @below {{op expected symbol reference @add_f32 to point to a reduction declaration}}
- omp.task in_reduction(@add_f32 -> %ptr : !llvm.ptr<f32>) {
+ omp.task in_reduction(@add_f32 -> %ptr : !llvm.ptr) {
// CHECK: "test.foo"() : () -> ()
"test.foo"() : () -> ()
// CHECK: omp.terminator
@@ -1324,9 +1324,9 @@ combiner {
omp.yield (%1 : f32)
}
-func.func @omp_task(%ptr: !llvm.ptr<f32>) {
+func.func @omp_task(%ptr: !llvm.ptr) {
// expected-error @below {{op accumulator variable used more than once}}
- omp.task in_reduction(@add_f32 -> %ptr : !llvm.ptr<f32>, @add_f32 -> %ptr : !llvm.ptr<f32>) {
+ omp.task in_reduction(@add_f32 -> %ptr : !llvm.ptr, @add_f32 -> %ptr : !llvm.ptr) {
// CHECK: "test.foo"() : () -> ()
"test.foo"() : () -> ()
// CHECK: omp.terminator
@@ -1348,14 +1348,14 @@ combiner {
omp.yield (%1 : i32)
}
atomic {
-^bb2(%arg2: !llvm.ptr<i32>, %arg3: !llvm.ptr<i32>):
- %2 = llvm.load %arg3 : !llvm.ptr<i32>
- llvm.atomicrmw add %arg2, %2 monotonic : !llvm.ptr<i32>, i32
+^bb2(%arg2: !llvm.ptr, %arg3: !llvm.ptr):
+ %2 = llvm.load %arg3 : !llvm.ptr -> i32
+ llvm.atomicrmw add %arg2, %2 monotonic : !llvm.ptr, i32
omp.yield
}
func.func @omp_task(%mem: memref<1xf32>) {
- // expected-error @below {{op expected accumulator ('memref<1xf32>') to be the same type as reduction declaration ('!llvm.ptr<i32>')}}
+ // expected-error @below {{op expected accumulator ('memref<1xf32>') to be the same type as reduction declaration ('!llvm.ptr')}}
omp.task in_reduction(@add_i32 -> %mem : memref<1xf32>) {
// CHECK: "test.foo"() : () -> ()
"test.foo"() : () -> ()
@@ -1493,52 +1493,52 @@ func.func @taskloop(%lb: i32, %ub: i32, %step: i32) {
// -----
func.func @taskloop(%lb: i32, %ub: i32, %step: i32) {
- %testf32 = "test.f32"() : () -> (!llvm.ptr<f32>)
- %testf32_2 = "test.f32"() : () -> (!llvm.ptr<f32>)
+ %testf32 = "test.f32"() : () -> (!llvm.ptr)
+ %testf32_2 = "test.f32"() : () -> (!llvm.ptr)
// expected-error @below {{expected as many reduction symbol references as reduction variables}}
"omp.taskloop"(%lb, %ub, %ub, %lb, %step, %step, %testf32, %testf32_2) ({
^bb0(%arg3: i32, %arg4: i32):
"omp.terminator"() : () -> ()
- }) {operandSegmentSizes = array<i32: 2, 2, 2, 0, 0, 0, 2, 0, 0, 0, 0, 0>, reductions = [@add_f32]} : (i32, i32, i32, i32, i32, i32, !llvm.ptr<f32>, !llvm.ptr<f32>) -> ()
+ }) {operandSegmentSizes = array<i32: 2, 2, 2, 0, 0, 0, 2, 0, 0, 0, 0, 0>, reductions = [@add_f32]} : (i32, i32, i32, i32, i32, i32, !llvm.ptr, !llvm.ptr) -> ()
return
}
// -----
func.func @taskloop(%lb: i32, %ub: i32, %step: i32) {
- %testf32 = "test.f32"() : () -> (!llvm.ptr<f32>)
- %testf32_2 = "test.f32"() : () -> (!llvm.ptr<f32>)
+ %testf32 = "test.f32"() : () -> (!llvm.ptr)
+ %testf32_2 = "test.f32"() : () -> (!llvm.ptr)
// expected-error @below {{expected as many reduction symbol references as reduction variables}}
"omp.taskloop"(%lb, %ub, %ub, %lb, %step, %step, %testf32) ({
^bb0(%arg3: i32, %arg4: i32):
"omp.terminator"() : () -> ()
- }) {operandSegmentSizes = array<i32: 2, 2, 2, 0, 0, 0, 1, 0, 0, 0, 0, 0>, reductions = [@add_f32, @add_f32]} : (i32, i32, i32, i32, i32, i32, !llvm.ptr<f32>) -> ()
+ }) {operandSegmentSizes = array<i32: 2, 2, 2, 0, 0, 0, 1, 0, 0, 0, 0, 0>, reductions = [@add_f32, @add_f32]} : (i32, i32, i32, i32, i32, i32, !llvm.ptr) -> ()
return
}
// -----
func.func @taskloop(%lb: i32, %ub: i32, %step: i32) {
- %testf32 = "test.f32"() : () -> (!llvm.ptr<f32>)
- %testf32_2 = "test.f32"() : () -> (!llvm.ptr<f32>)
+ %testf32 = "test.f32"() : () -> (!llvm.ptr)
+ %testf32_2 = "test.f32"() : () -> (!llvm.ptr)
// expected-error @below {{expected as many reduction symbol references as reduction variables}}
"omp.taskloop"(%lb, %ub, %ub, %lb, %step, %step, %testf32, %testf32_2) ({
^bb0(%arg3: i32, %arg4: i32):
"omp.terminator"() : () -> ()
- }) {in_reductions = [@add_f32], operandSegmentSizes = array<i32: 2, 2, 2, 0, 0, 2, 0, 0, 0, 0, 0, 0>} : (i32, i32, i32, i32, i32, i32, !llvm.ptr<f32>, !llvm.ptr<f32>) -> ()
+ }) {in_reductions = [@add_f32], operandSegmentSizes = array<i32: 2, 2, 2, 0, 0, 2, 0, 0, 0, 0, 0, 0>} : (i32, i32, i32, i32, i32, i32, !llvm.ptr, !llvm.ptr) -> ()
return
}
// -----
func.func @taskloop(%lb: i32, %ub: i32, %step: i32) {
- %testf32 = "test.f32"() : () -> (!llvm.ptr<f32>)
- %testf32_2 = "test.f32"() : () -> (!llvm.ptr<f32>)
+ %testf32 = "test.f32"() : () -> (!llvm.ptr)
+ %testf32_2 = "test.f32"() : () -> (!llvm.ptr)
// expected-error @below {{expected as many reduction symbol references as reduction variables}}
"omp.taskloop"(%lb, %ub, %ub, %lb, %step, %step, %testf32_2) ({
^bb0(%arg3: i32, %arg4: i32):
"omp.terminator"() : () -> ()
- }) {in_reductions = [@add_f32, @add_f32], operandSegmentSizes = array<i32: 2, 2, 2, 0, 0, 1, 0, 0, 0, 0, 0, 0>} : (i32, i32, i32, i32, i32, i32, !llvm.ptr<f32>) -> ()
+ }) {in_reductions = [@add_f32, @add_f32], operandSegmentSizes = array<i32: 2, 2, 2, 0, 0, 1, 0, 0, 0, 0, 0, 0>} : (i32, i32, i32, i32, i32, i32, !llvm.ptr) -> ()
return
}
@@ -1557,10 +1557,10 @@ combiner {
}
func.func @taskloop(%lb: i32, %ub: i32, %step: i32) {
- %testf32 = "test.f32"() : () -> (!llvm.ptr<f32>)
- %testf32_2 = "test.f32"() : () -> (!llvm.ptr<f32>)
+ %testf32 = "test.f32"() : () -> (!llvm.ptr)
+ %testf32_2 = "test.f32"() : () -> (!llvm.ptr)
// expected-error @below {{if a reduction clause is present on the taskloop directive, the nogroup clause must not be specified}}
- omp.taskloop reduction(@add_f32 -> %testf32 : !llvm.ptr<f32>, @add_f32 -> %testf32_2 : !llvm.ptr<f32>) nogroup
+ omp.taskloop reduction(@add_f32 -> %testf32 : !llvm.ptr, @add_f32 -> %testf32_2 : !llvm.ptr) nogroup
for (%i, %j) : i32 = (%lb, %ub) to (%ub, %lb) step (%step, %step) {
omp.terminator
}
@@ -1582,9 +1582,9 @@ combiner {
}
func.func @taskloop(%lb: i32, %ub: i32, %step: i32) {
- %testf32 = "test.f32"() : () -> (!llvm.ptr<f32>)
+ %testf32 = "test.f32"() : () -> (!llvm.ptr)
// expected-error @below {{the same list item cannot appear in both a reduction and an in_reduction clause}}
- omp.taskloop reduction(@add_f32 -> %testf32 : !llvm.ptr<f32>) in_reduction(@add_f32 -> %testf32 : !llvm.ptr<f32>)
+ omp.taskloop reduction(@add_f32 -> %testf32 : !llvm.ptr) in_reduction(@add_f32 -> %testf32 : !llvm.ptr)
for (%i, %j) : i32 = (%lb, %ub) to (%ub, %lb) step (%step, %step) {
omp.terminator
}
@@ -1606,9 +1606,9 @@ func.func @taskloop(%lb: i32, %ub: i32, %step: i32) {
// -----
func.func @omp_threadprivate() {
- %1 = llvm.mlir.addressof @_QFsubEx : !llvm.ptr<i32>
+ %1 = llvm.mlir.addressof @_QFsubEx : !llvm.ptr
// expected-error @below {{op failed to verify that all of {sym_addr, tls_addr} have same type}}
- %2 = omp.threadprivate %1 : !llvm.ptr<i32> -> memref<i32>
+ %2 = omp.threadprivate %1 : !llvm.ptr -> memref<i32>
return
}
diff --git a/mlir/test/Dialect/OpenMP/ops.mlir b/mlir/test/Dialect/OpenMP/ops.mlir
index 47788be6a7b2c9b..d59a4f428118bf3 100644
--- a/mlir/test/Dialect/OpenMP/ops.mlir
+++ b/mlir/test/Dialect/OpenMP/ops.mlir
@@ -565,7 +565,7 @@ func.func @omp_target_pretty(%if_cond : i1, %device : si32, %num_threads : i32)
// CHECK: ^{{.+}}(%{{.+}}: f32, %{{.+}}: f32):
// CHECK: omp.yield
// CHECK: atomic
-// CHECK: ^{{.+}}(%{{.+}}: !llvm.ptr<f32>, %{{.+}}: !llvm.ptr<f32>):
+// CHECK: ^{{.+}}(%{{.+}}: !llvm.ptr, %{{.+}}: !llvm.ptr):
// CHECK: omp.yield
omp.reduction.declare @add_f32 : f32
init {
@@ -579,22 +579,22 @@ combiner {
omp.yield (%1 : f32)
}
atomic {
-^bb2(%arg2: !llvm.ptr<f32>, %arg3: !llvm.ptr<f32>):
- %2 = llvm.load %arg3 : !llvm.ptr<f32>
- llvm.atomicrmw fadd %arg2, %2 monotonic : !llvm.ptr<f32>, f32
+^bb2(%arg2: !llvm.ptr, %arg3: !llvm.ptr):
+ %2 = llvm.load %arg3 : !llvm.ptr -> f32
+ llvm.atomicrmw fadd %arg2, %2 monotonic : !llvm.ptr, f32
omp.yield
}
// CHECK-LABEL: func @wsloop_reduction
func.func @wsloop_reduction(%lb : index, %ub : index, %step : index) {
%c1 = arith.constant 1 : i32
- %0 = llvm.alloca %c1 x i32 : (i32) -> !llvm.ptr<f32>
- // CHECK: reduction(@add_f32 -> %{{.+}} : !llvm.ptr<f32>)
- omp.wsloop reduction(@add_f32 -> %0 : !llvm.ptr<f32>)
+ %0 = llvm.alloca %c1 x i32 : (i32) -> !llvm.ptr
+ // CHECK: reduction(@add_f32 -> %{{.+}} : !llvm.ptr)
+ omp.wsloop reduction(@add_f32 -> %0 : !llvm.ptr)
for (%iv) : index = (%lb) to (%ub) step (%step) {
%1 = arith.constant 2.0 : f32
// CHECK: omp.reduction %{{.+}}, %{{.+}}
- omp.reduction %1, %0 : f32, !llvm.ptr<f32>
+ omp.reduction %1, %0 : f32, !llvm.ptr
omp.yield
}
return
@@ -603,12 +603,12 @@ func.func @wsloop_reduction(%lb : index, %ub : index, %step : index) {
// CHECK-LABEL: func @parallel_reduction
func.func @parallel_reduction() {
%c1 = arith.constant 1 : i32
- %0 = llvm.alloca %c1 x i32 : (i32) -> !llvm.ptr<f32>
- // CHECK: omp.parallel reduction(@add_f32 -> {{.+}} : !llvm.ptr<f32>)
- omp.parallel reduction(@add_f32 -> %0 : !llvm.ptr<f32>) {
+ %0 = llvm.alloca %c1 x i32 : (i32) -> !llvm.ptr
+ // CHECK: omp.parallel reduction(@add_f32 -> {{.+}} : !llvm.ptr)
+ omp.parallel reduction(@add_f32 -> %0 : !llvm.ptr) {
%1 = arith.constant 2.0 : f32
// CHECK: omp.reduction %{{.+}}, %{{.+}}
- omp.reduction %1, %0 : f32, !llvm.ptr<f32>
+ omp.reduction %1, %0 : f32, !llvm.ptr
omp.terminator
}
return
@@ -617,14 +617,14 @@ func.func @parallel_reduction() {
// CHECK: func @parallel_wsloop_reduction
func.func @parallel_wsloop_reduction(%lb : index, %ub : index, %step : index) {
%c1 = arith.constant 1 : i32
- %0 = llvm.alloca %c1 x i32 : (i32) -> !llvm.ptr<f32>
- // CHECK: omp.parallel reduction(@add_f32 -> %{{.+}} : !llvm.ptr<f32>) {
- omp.parallel reduction(@add_f32 -> %0 : !llvm.ptr<f32>) {
+ %0 = llvm.alloca %c1 x i32 : (i32) -> !llvm.ptr
+ // CHECK: omp.parallel reduction(@add_f32 -> %{{.+}} : !llvm.ptr) {
+ omp.parallel reduction(@add_f32 -> %0 : !llvm.ptr) {
// CHECK: omp.wsloop for (%{{.+}}) : index = (%{{.+}}) to (%{{.+}}) step (%{{.+}})
omp.wsloop for (%iv) : index = (%lb) to (%ub) step (%step) {
%1 = arith.constant 2.0 : f32
- // CHECK: omp.reduction %{{.+}}, %{{.+}} : f32, !llvm.ptr<f32>
- omp.reduction %1, %0 : f32, !llvm.ptr<f32>
+ // CHECK: omp.reduction %{{.+}}, %{{.+}} : f32, !llvm.ptr
+ omp.reduction %1, %0 : f32, !llvm.ptr
// CHECK: omp.yield
omp.yield
}
@@ -684,12 +684,12 @@ func.func @omp_teams(%lb : i32, %ub : i32, %if_cond : i1, %num_threads : i32,
// Test reduction.
%c1 = arith.constant 1 : i32
- %0 = llvm.alloca %c1 x i32 : (i32) -> !llvm.ptr<f32>
- // CHECK: omp.teams reduction(@add_f32 -> %{{.+}} : !llvm.ptr<f32>) {
- omp.teams reduction(@add_f32 -> %0 : !llvm.ptr<f32>) {
+ %0 = llvm.alloca %c1 x i32 : (i32) -> !llvm.ptr
+ // CHECK: omp.teams reduction(@add_f32 -> %{{.+}} : !llvm.ptr) {
+ omp.teams reduction(@add_f32 -> %0 : !llvm.ptr) {
%1 = arith.constant 2.0 : f32
// CHECK: omp.reduction %{{.+}}, %{{.+}}
- omp.reduction %1, %0 : f32, !llvm.ptr<f32>
+ omp.reduction %1, %0 : f32, !llvm.ptr
// CHECK: omp.terminator
omp.terminator
}
@@ -707,21 +707,21 @@ func.func @omp_teams(%lb : i32, %ub : i32, %if_cond : i1, %num_threads : i32,
// CHECK-LABEL: func @sections_reduction
func.func @sections_reduction() {
%c1 = arith.constant 1 : i32
- %0 = llvm.alloca %c1 x i32 : (i32) -> !llvm.ptr<f32>
- // CHECK: omp.sections reduction(@add_f32 -> {{.+}} : !llvm.ptr<f32>)
- omp.sections reduction(@add_f32 -> %0 : !llvm.ptr<f32>) {
+ %0 = llvm.alloca %c1 x i32 : (i32) -> !llvm.ptr
+ // CHECK: omp.sections reduction(@add_f32 -> {{.+}} : !llvm.ptr)
+ omp.sections reduction(@add_f32 -> %0 : !llvm.ptr) {
// CHECK: omp.section
omp.section {
%1 = arith.constant 2.0 : f32
// CHECK: omp.reduction %{{.+}}, %{{.+}}
- omp.reduction %1, %0 : f32, !llvm.ptr<f32>
+ omp.reduction %1, %0 : f32, !llvm.ptr
omp.terminator
}
// CHECK: omp.section
omp.section {
%1 = arith.constant 3.0 : f32
// CHECK: omp.reduction %{{.+}}, %{{.+}}
- omp.reduction %1, %0 : f32, !llvm.ptr<f32>
+ omp.reduction %1, %0 : f32, !llvm.ptr
omp.terminator
}
omp.terminator
@@ -776,14 +776,14 @@ func.func @parallel_reduction2() {
// CHECK: func @parallel_wsloop_reduction2
func.func @parallel_wsloop_reduction2(%lb : index, %ub : index, %step : index) {
%c1 = arith.constant 1 : i32
- %0 = llvm.alloca %c1 x i32 : (i32) -> !llvm.ptr<f32>
- // CHECK: omp.parallel reduction(@add2_f32 -> %{{.+}} : !llvm.ptr<f32>) {
- omp.parallel reduction(@add2_f32 -> %0 : !llvm.ptr<f32>) {
+ %0 = llvm.alloca %c1 x i32 : (i32) -> !llvm.ptr
+ // CHECK: omp.parallel reduction(@add2_f32 -> %{{.+}} : !llvm.ptr) {
+ omp.parallel reduction(@add2_f32 -> %0 : !llvm.ptr) {
// CHECK: omp.wsloop for (%{{.+}}) : index = (%{{.+}}) to (%{{.+}}) step (%{{.+}})
omp.wsloop for (%iv) : index = (%lb) to (%ub) step (%step) {
%1 = arith.constant 2.0 : f32
- // CHECK: omp.reduction %{{.+}}, %{{.+}} : f32, !llvm.ptr<f32>
- omp.reduction %1, %0 : f32, !llvm.ptr<f32>
+ // CHECK: omp.reduction %{{.+}}, %{{.+}} : f32, !llvm.ptr
+ omp.reduction %1, %0 : f32, !llvm.ptr
// CHECK: omp.yield
omp.yield
}
@@ -1441,18 +1441,18 @@ func.func @omp_atomic_capture(%v: memref<i32>, %x: memref<i32>, %expr: i32) {
// CHECK-LABEL: omp_sectionsop
func.func @omp_sectionsop(%data_var1 : memref<i32>, %data_var2 : memref<i32>,
- %data_var3 : memref<i32>, %redn_var : !llvm.ptr<f32>) {
+ %data_var3 : memref<i32>, %redn_var : !llvm.ptr) {
// CHECK: omp.sections allocate(%{{.*}} : memref<i32> -> %{{.*}} : memref<i32>)
"omp.sections" (%data_var1, %data_var1) ({
// CHECK: omp.terminator
omp.terminator
}) {operandSegmentSizes = array<i32: 0,1,1>} : (memref<i32>, memref<i32>) -> ()
- // CHECK: omp.sections reduction(@add_f32 -> %{{.*}} : !llvm.ptr<f32>)
+ // CHECK: omp.sections reduction(@add_f32 -> %{{.*}} : !llvm.ptr)
"omp.sections" (%redn_var) ({
// CHECK: omp.terminator
omp.terminator
- }) {operandSegmentSizes = array<i32: 1,0,0>, reductions=[@add_f32]} : (!llvm.ptr<f32>) -> ()
+ }) {operandSegmentSizes = array<i32: 1,0,0>, reductions=[@add_f32]} : (!llvm.ptr) -> ()
// CHECK: omp.sections nowait {
omp.sections nowait {
@@ -1460,8 +1460,8 @@ func.func @omp_sectionsop(%data_var1 : memref<i32>, %data_var2 : memref<i32>,
omp.terminator
}
- // CHECK: omp.sections reduction(@add_f32 -> %{{.*}} : !llvm.ptr<f32>) {
- omp.sections reduction(@add_f32 -> %redn_var : !llvm.ptr<f32>) {
+ // CHECK: omp.sections reduction(@add_f32 -> %{{.*}} : !llvm.ptr) {
+ omp.sections reduction(@add_f32 -> %redn_var : !llvm.ptr) {
// CHECK: omp.terminator
omp.terminator
}
@@ -1485,13 +1485,13 @@ func.func @omp_sectionsop(%data_var1 : memref<i32>, %data_var2 : memref<i32>,
}
// CHECK: omp.section
omp.section {
- // CHECK: %{{.*}} = "test.payload"(%{{.*}}) : (!llvm.ptr<f32>) -> i32
- %1 = "test.payload"(%redn_var) : (!llvm.ptr<f32>) -> i32
+ // CHECK: %{{.*}} = "test.payload"(%{{.*}}) : (!llvm.ptr) -> i32
+ %1 = "test.payload"(%redn_var) : (!llvm.ptr) -> i32
}
// CHECK: omp.section
omp.section {
- // CHECK: "test.payload"(%{{.*}}) : (!llvm.ptr<f32>) -> ()
- "test.payload"(%redn_var) : (!llvm.ptr<f32>) -> ()
+ // CHECK: "test.payload"(%{{.*}}) : (!llvm.ptr) -> ()
+ "test.payload"(%redn_var) : (!llvm.ptr) -> ()
}
// CHECK: omp.terminator
omp.terminator
@@ -1613,12 +1613,12 @@ func.func @omp_task(%bool_var: i1, %i64_var: i64, %i32_var: i32, %data_var: memr
// Checking `in_reduction` clause
%c1 = arith.constant 1 : i32
- // CHECK: %[[redn_var1:.*]] = llvm.alloca %{{.*}} x f32 : (i32) -> !llvm.ptr<f32>
- %0 = llvm.alloca %c1 x f32 : (i32) -> !llvm.ptr<f32>
- // CHECK: %[[redn_var2:.*]] = llvm.alloca %{{.*}} x f32 : (i32) -> !llvm.ptr<f32>
- %1 = llvm.alloca %c1 x f32 : (i32) -> !llvm.ptr<f32>
- // CHECK: omp.task in_reduction(@add_f32 -> %[[redn_var1]] : !llvm.ptr<f32>, @add_f32 -> %[[redn_var2]] : !llvm.ptr<f32>) {
- omp.task in_reduction(@add_f32 -> %0 : !llvm.ptr<f32>, @add_f32 -> %1 : !llvm.ptr<f32>) {
+ // CHECK: %[[redn_var1:.*]] = llvm.alloca %{{.*}} x f32 : (i32) -> !llvm.ptr
+ %0 = llvm.alloca %c1 x f32 : (i32) -> !llvm.ptr
+ // CHECK: %[[redn_var2:.*]] = llvm.alloca %{{.*}} x f32 : (i32) -> !llvm.ptr
+ %1 = llvm.alloca %c1 x f32 : (i32) -> !llvm.ptr
+ // CHECK: omp.task in_reduction(@add_f32 -> %[[redn_var1]] : !llvm.ptr, @add_f32 -> %[[redn_var2]] : !llvm.ptr) {
+ omp.task in_reduction(@add_f32 -> %0 : !llvm.ptr, @add_f32 -> %1 : !llvm.ptr) {
// CHECK: "test.foo"() : () -> ()
"test.foo"() : () -> ()
// CHECK: omp.terminator
@@ -1646,8 +1646,8 @@ func.func @omp_task(%bool_var: i1, %i64_var: i64, %i32_var: i32, %data_var: memr
// Checking multiple clauses
// CHECK: omp.task if(%[[bool_var]]) final(%[[bool_var]]) untied
omp.task if(%bool_var) final(%bool_var) untied
- // CHECK-SAME: in_reduction(@add_f32 -> %[[redn_var1]] : !llvm.ptr<f32>, @add_f32 -> %[[redn_var2]] : !llvm.ptr<f32>)
- in_reduction(@add_f32 -> %0 : !llvm.ptr<f32>, @add_f32 -> %1 : !llvm.ptr<f32>)
+ // CHECK-SAME: in_reduction(@add_f32 -> %[[redn_var1]] : !llvm.ptr, @add_f32 -> %[[redn_var2]] : !llvm.ptr)
+ in_reduction(@add_f32 -> %0 : !llvm.ptr, @add_f32 -> %1 : !llvm.ptr)
// CHECK-SAME: priority(%[[i32_var]])
priority(%i32_var)
// CHECK-SAME: allocate(%[[data_var]] : memref<i32> -> %[[data_var]] : memref<i32>)
@@ -1679,20 +1679,20 @@ func.func @omp_threadprivate() {
%1 = arith.constant 2 : i32
%2 = arith.constant 3 : i32
- // CHECK: [[ARG0:%.*]] = llvm.mlir.addressof @_QFsubEx : !llvm.ptr<i32>
- // CHECK: {{.*}} = omp.threadprivate [[ARG0]] : !llvm.ptr<i32> -> !llvm.ptr<i32>
- %3 = llvm.mlir.addressof @_QFsubEx : !llvm.ptr<i32>
- %4 = omp.threadprivate %3 : !llvm.ptr<i32> -> !llvm.ptr<i32>
- llvm.store %0, %4 : !llvm.ptr<i32>
+ // CHECK: [[ARG0:%.*]] = llvm.mlir.addressof @_QFsubEx : !llvm.ptr
+ // CHECK: {{.*}} = omp.threadprivate [[ARG0]] : !llvm.ptr -> !llvm.ptr
+ %3 = llvm.mlir.addressof @_QFsubEx : !llvm.ptr
+ %4 = omp.threadprivate %3 : !llvm.ptr -> !llvm.ptr
+ llvm.store %0, %4 : i32, !llvm.ptr
// CHECK: omp.parallel
- // CHECK: {{.*}} = omp.threadprivate [[ARG0]] : !llvm.ptr<i32> -> !llvm.ptr<i32>
+ // CHECK: {{.*}} = omp.threadprivate [[ARG0]] : !llvm.ptr -> !llvm.ptr
omp.parallel {
- %5 = omp.threadprivate %3 : !llvm.ptr<i32> -> !llvm.ptr<i32>
- llvm.store %1, %5 : !llvm.ptr<i32>
+ %5 = omp.threadprivate %3 : !llvm.ptr -> !llvm.ptr
+ llvm.store %1, %5 : i32, !llvm.ptr
omp.terminator
}
- llvm.store %2, %4 : !llvm.ptr<i32>
+ llvm.store %2, %4 : i32, !llvm.ptr
return
}
@@ -1810,9 +1810,9 @@ func.func @omp_taskgroup_multiple_tasks() -> () {
// CHECK-LABEL: @omp_taskgroup_clauses
func.func @omp_taskgroup_clauses() -> () {
%testmemref = "test.memref"() : () -> (memref<i32>)
- %testf32 = "test.f32"() : () -> (!llvm.ptr<f32>)
- // CHECK: omp.taskgroup task_reduction(@add_f32 -> %{{.+}}: !llvm.ptr<f32>) allocate(%{{.+}}: memref<i32> -> %{{.+}}: memref<i32>)
- omp.taskgroup allocate(%testmemref : memref<i32> -> %testmemref : memref<i32>) task_reduction(@add_f32 -> %testf32 : !llvm.ptr<f32>) {
+ %testf32 = "test.f32"() : () -> (!llvm.ptr)
+ // CHECK: omp.taskgroup task_reduction(@add_f32 -> %{{.+}}: !llvm.ptr) allocate(%{{.+}}: memref<i32> -> %{{.+}}: memref<i32>)
+ omp.taskgroup allocate(%testmemref : memref<i32> -> %testmemref : memref<i32>) task_reduction(@add_f32 -> %testf32 : !llvm.ptr) {
// CHECK: omp.task
omp.task {
"test.foo"() : () -> ()
@@ -1896,27 +1896,27 @@ func.func @omp_taskloop(%lb: i32, %ub: i32, %step: i32) -> () {
omp.terminator
}
- %testf32 = "test.f32"() : () -> (!llvm.ptr<f32>)
- %testf32_2 = "test.f32"() : () -> (!llvm.ptr<f32>)
- // CHECK: omp.taskloop in_reduction(@add_f32 -> %{{.+}} : !llvm.ptr<f32>, @add_f32 -> %{{.+}} : !llvm.ptr<f32>)
+ %testf32 = "test.f32"() : () -> (!llvm.ptr)
+ %testf32_2 = "test.f32"() : () -> (!llvm.ptr)
+ // CHECK: omp.taskloop in_reduction(@add_f32 -> %{{.+}} : !llvm.ptr, @add_f32 -> %{{.+}} : !llvm.ptr)
// CHECK-SAME: for (%{{.+}}, %{{.+}}) : i32 = (%{{.+}}, %{{.+}}) to (%{{.+}}, %{{.+}}) step (%{{.+}}, %{{.+}}) {
- omp.taskloop in_reduction(@add_f32 -> %testf32 : !llvm.ptr<f32>, @add_f32 -> %testf32_2 : !llvm.ptr<f32>)
+ omp.taskloop in_reduction(@add_f32 -> %testf32 : !llvm.ptr, @add_f32 -> %testf32_2 : !llvm.ptr)
for (%i, %j) : i32 = (%lb, %ub) to (%ub, %lb) step (%step, %step) {
// CHECK: omp.terminator
omp.terminator
}
- // CHECK: omp.taskloop reduction(@add_f32 -> %{{.+}} : !llvm.ptr<f32>, @add_f32 -> %{{.+}} : !llvm.ptr<f32>)
+ // CHECK: omp.taskloop reduction(@add_f32 -> %{{.+}} : !llvm.ptr, @add_f32 -> %{{.+}} : !llvm.ptr)
// CHECK-SAME: for (%{{.+}}, %{{.+}}) : i32 = (%{{.+}}, %{{.+}}) to (%{{.+}}, %{{.+}}) step (%{{.+}}, %{{.+}}) {
- omp.taskloop reduction(@add_f32 -> %testf32 : !llvm.ptr<f32>, @add_f32 -> %testf32_2 : !llvm.ptr<f32>)
+ omp.taskloop reduction(@add_f32 -> %testf32 : !llvm.ptr, @add_f32 -> %testf32_2 : !llvm.ptr)
for (%i, %j) : i32 = (%lb, %ub) to (%ub, %lb) step (%step, %step) {
// CHECK: omp.terminator
omp.terminator
}
- // CHECK: omp.taskloop in_reduction(@add_f32 -> %{{.+}} : !llvm.ptr<f32>) reduction(@add_f32 -> %{{.+}} : !llvm.ptr<f32>)
+ // CHECK: omp.taskloop in_reduction(@add_f32 -> %{{.+}} : !llvm.ptr) reduction(@add_f32 -> %{{.+}} : !llvm.ptr)
// CHECK-SAME: for (%{{.+}}, %{{.+}}) : i32 = (%{{.+}}, %{{.+}}) to (%{{.+}}, %{{.+}}) step (%{{.+}}, %{{.+}}) {
- omp.taskloop in_reduction(@add_f32 -> %testf32 : !llvm.ptr<f32>) reduction(@add_f32 -> %testf32_2 : !llvm.ptr<f32>)
+ omp.taskloop in_reduction(@add_f32 -> %testf32 : !llvm.ptr) reduction(@add_f32 -> %testf32_2 : !llvm.ptr)
for (%i, %j) : i32 = (%lb, %ub) to (%ub, %lb) step (%step, %step) {
// CHECK: omp.terminator
omp.terminator
@@ -2027,49 +2027,49 @@ atomic {
}
// CHECK-LABEL: omp_targets_with_map_bounds
-// CHECK-SAME: (%[[ARG0:.*]]: !llvm.ptr<array<10 x i32>>, %[[ARG1:.*]]: !llvm.ptr<array<10 x i32>>)
-func.func @omp_targets_with_map_bounds(%arg0: !llvm.ptr<array<10 x i32>>, %arg1: !llvm.ptr<array<10 x i32>>) -> () {
+// CHECK-SAME: (%[[ARG0:.*]]: !llvm.ptr, %[[ARG1:.*]]: !llvm.ptr)
+func.func @omp_targets_with_map_bounds(%arg0: !llvm.ptr, %arg1: !llvm.ptr) -> () {
// CHECK: %[[C_00:.*]] = llvm.mlir.constant(4 : index) : i64
// CHECK: %[[C_01:.*]] = llvm.mlir.constant(1 : index) : i64
// CHECK: %[[C_02:.*]] = llvm.mlir.constant(1 : index) : i64
// CHECK: %[[C_03:.*]] = llvm.mlir.constant(1 : index) : i64
// CHECK: %[[BOUNDS0:.*]] = omp.bounds lower_bound(%[[C_01]] : i64) upper_bound(%[[C_00]] : i64) stride(%[[C_02]] : i64) start_idx(%[[C_03]] : i64)
- // CHECK: %[[MAP0:.*]] = omp.map_info var_ptr(%[[ARG0]] : !llvm.ptr<array<10 x i32>>, !llvm.array<10 x i32>) map_clauses(tofrom) capture(ByRef) bounds(%[[BOUNDS0]]) -> !llvm.ptr<array<10 x i32>> {name = ""}
+ // CHECK: %[[MAP0:.*]] = omp.map_info var_ptr(%[[ARG0]] : !llvm.ptr, !llvm.array<10 x i32>) map_clauses(tofrom) capture(ByRef) bounds(%[[BOUNDS0]]) -> !llvm.ptr {name = ""}
%0 = llvm.mlir.constant(4 : index) : i64
%1 = llvm.mlir.constant(1 : index) : i64
%2 = llvm.mlir.constant(1 : index) : i64
%3 = llvm.mlir.constant(1 : index) : i64
%4 = omp.bounds lower_bound(%1 : i64) upper_bound(%0 : i64) stride(%2 : i64) start_idx(%3 : i64)
- %mapv1 = omp.map_info var_ptr(%arg0 : !llvm.ptr<array<10 x i32>>, !llvm.array<10 x i32>) map_clauses(tofrom) capture(ByRef) bounds(%4) -> !llvm.ptr<array<10 x i32>> {name = ""}
+ %mapv1 = omp.map_info var_ptr(%arg0 : !llvm.ptr, !llvm.array<10 x i32>) map_clauses(tofrom) capture(ByRef) bounds(%4) -> !llvm.ptr {name = ""}
// CHECK: %[[C_10:.*]] = llvm.mlir.constant(9 : index) : i64
// CHECK: %[[C_11:.*]] = llvm.mlir.constant(1 : index) : i64
// CHECK: %[[C_12:.*]] = llvm.mlir.constant(2 : index) : i64
// CHECK: %[[C_13:.*]] = llvm.mlir.constant(2 : index) : i64
// CHECK: %[[BOUNDS1:.*]] = omp.bounds lower_bound(%[[C_11]] : i64) upper_bound(%[[C_10]] : i64) stride(%[[C_12]] : i64) start_idx(%[[C_13]] : i64)
- // CHECK: %[[MAP1:.*]] = omp.map_info var_ptr(%[[ARG1]] : !llvm.ptr<array<10 x i32>>, !llvm.array<10 x i32>) map_clauses(exit_release_or_enter_alloc) capture(ByCopy) bounds(%[[BOUNDS1]]) -> !llvm.ptr<array<10 x i32>> {name = ""}
+ // CHECK: %[[MAP1:.*]] = omp.map_info var_ptr(%[[ARG1]] : !llvm.ptr, !llvm.array<10 x i32>) map_clauses(exit_release_or_enter_alloc) capture(ByCopy) bounds(%[[BOUNDS1]]) -> !llvm.ptr {name = ""}
%6 = llvm.mlir.constant(9 : index) : i64
%7 = llvm.mlir.constant(1 : index) : i64
%8 = llvm.mlir.constant(2 : index) : i64
%9 = llvm.mlir.constant(2 : index) : i64
%10 = omp.bounds lower_bound(%7 : i64) upper_bound(%6 : i64) stride(%8 : i64) start_idx(%9 : i64)
- %mapv2 = omp.map_info var_ptr(%arg1 : !llvm.ptr<array<10 x i32>>, !llvm.array<10 x i32>) map_clauses(exit_release_or_enter_alloc) capture(ByCopy) bounds(%10) -> !llvm.ptr<array<10 x i32>> {name = ""}
+ %mapv2 = omp.map_info var_ptr(%arg1 : !llvm.ptr, !llvm.array<10 x i32>) map_clauses(exit_release_or_enter_alloc) capture(ByCopy) bounds(%10) -> !llvm.ptr {name = ""}
- // CHECK: omp.target map_entries(%[[MAP0]], %[[MAP1]] : !llvm.ptr<array<10 x i32>>, !llvm.ptr<array<10 x i32>>)
- omp.target map_entries(%mapv1, %mapv2 : !llvm.ptr<array<10 x i32>>, !llvm.ptr<array<10 x i32>>){}
+ // CHECK: omp.target map_entries(%[[MAP0]], %[[MAP1]] : !llvm.ptr, !llvm.ptr)
+ omp.target map_entries(%mapv1, %mapv2 : !llvm.ptr, !llvm.ptr){}
- // CHECK: omp.target_data map_entries(%[[MAP0]], %[[MAP1]] : !llvm.ptr<array<10 x i32>>, !llvm.ptr<array<10 x i32>>)
- omp.target_data map_entries(%mapv1, %mapv2 : !llvm.ptr<array<10 x i32>>, !llvm.ptr<array<10 x i32>>){}
+ // CHECK: omp.target_data map_entries(%[[MAP0]], %[[MAP1]] : !llvm.ptr, !llvm.ptr)
+ omp.target_data map_entries(%mapv1, %mapv2 : !llvm.ptr, !llvm.ptr){}
- // CHECK: %[[MAP2:.*]] = omp.map_info var_ptr(%[[ARG0]] : !llvm.ptr<array<10 x i32>>, !llvm.array<10 x i32>) map_clauses(exit_release_or_enter_alloc) capture(VLAType) bounds(%[[BOUNDS0]]) -> !llvm.ptr<array<10 x i32>> {name = ""}
- // CHECK: omp.target_enter_data map_entries(%[[MAP2]] : !llvm.ptr<array<10 x i32>>)
- %mapv3 = omp.map_info var_ptr(%arg0 : !llvm.ptr<array<10 x i32>>, !llvm.array<10 x i32>) map_clauses(exit_release_or_enter_alloc) capture(VLAType) bounds(%4) -> !llvm.ptr<array<10 x i32>> {name = ""}
- omp.target_enter_data map_entries(%mapv3 : !llvm.ptr<array<10 x i32>>){}
+ // CHECK: %[[MAP2:.*]] = omp.map_info var_ptr(%[[ARG0]] : !llvm.ptr, !llvm.array<10 x i32>) map_clauses(exit_release_or_enter_alloc) capture(VLAType) bounds(%[[BOUNDS0]]) -> !llvm.ptr {name = ""}
+ // CHECK: omp.target_enter_data map_entries(%[[MAP2]] : !llvm.ptr)
+ %mapv3 = omp.map_info var_ptr(%arg0 : !llvm.ptr, !llvm.array<10 x i32>) map_clauses(exit_release_or_enter_alloc) capture(VLAType) bounds(%4) -> !llvm.ptr {name = ""}
+ omp.target_enter_data map_entries(%mapv3 : !llvm.ptr){}
- // CHECK: %[[MAP3:.*]] = omp.map_info var_ptr(%[[ARG1]] : !llvm.ptr<array<10 x i32>>, !llvm.array<10 x i32>) map_clauses(exit_release_or_enter_alloc) capture(This) bounds(%[[BOUNDS1]]) -> !llvm.ptr<array<10 x i32>> {name = ""}
- // CHECK: omp.target_exit_data map_entries(%[[MAP3]] : !llvm.ptr<array<10 x i32>>)
- %mapv4 = omp.map_info var_ptr(%arg1 : !llvm.ptr<array<10 x i32>>, !llvm.array<10 x i32>) map_clauses(exit_release_or_enter_alloc) capture(This) bounds(%10) -> !llvm.ptr<array<10 x i32>> {name = ""}
- omp.target_exit_data map_entries(%mapv4 : !llvm.ptr<array<10 x i32>>){}
+ // CHECK: %[[MAP3:.*]] = omp.map_info var_ptr(%[[ARG1]] : !llvm.ptr, !llvm.array<10 x i32>) map_clauses(exit_release_or_enter_alloc) capture(This) bounds(%[[BOUNDS1]]) -> !llvm.ptr {name = ""}
+ // CHECK: omp.target_exit_data map_entries(%[[MAP3]] : !llvm.ptr)
+ %mapv4 = omp.map_info var_ptr(%arg1 : !llvm.ptr, !llvm.array<10 x i32>) map_clauses(exit_release_or_enter_alloc) capture(This) bounds(%10) -> !llvm.ptr {name = ""}
+ omp.target_exit_data map_entries(%mapv4 : !llvm.ptr){}
return
}
diff --git a/mlir/test/Dialect/SparseTensor/invalid.mlir b/mlir/test/Dialect/SparseTensor/invalid.mlir
index dd2298d4c3c143c..f85bc5111d7a27c 100644
--- a/mlir/test/Dialect/SparseTensor/invalid.mlir
+++ b/mlir/test/Dialect/SparseTensor/invalid.mlir
@@ -1,8 +1,8 @@
// RUN: mlir-opt %s -split-input-file -verify-diagnostics
-func.func @invalid_new_dense(%arg0: !llvm.ptr<i8>) -> tensor<32xf32> {
+func.func @invalid_new_dense(%arg0: !llvm.ptr) -> tensor<32xf32> {
// expected-error at +1 {{'sparse_tensor.new' op result #0 must be sparse tensor of any type values, but got 'tensor<32xf32>'}}
- %0 = sparse_tensor.new %arg0 : !llvm.ptr<i8> to tensor<32xf32>
+ %0 = sparse_tensor.new %arg0 : !llvm.ptr to tensor<32xf32>
return %0 : tensor<32xf32>
}
@@ -389,9 +389,9 @@ func.func @sparse_convert_dim_mismatch(%arg0: tensor<10x?xf32>) -> tensor<10x10x
// -----
-func.func @invalid_out_dense(%arg0: tensor<10xf64>, %arg1: !llvm.ptr<i8>) {
+func.func @invalid_out_dense(%arg0: tensor<10xf64>, %arg1: !llvm.ptr) {
// expected-error at +1 {{'sparse_tensor.out' op operand #0 must be sparse tensor of any type values, but got 'tensor<10xf64>'}}
- sparse_tensor.out %arg0, %arg1 : tensor<10xf64>, !llvm.ptr<i8>
+ sparse_tensor.out %arg0, %arg1 : tensor<10xf64>, !llvm.ptr
return
}
diff --git a/mlir/test/Dialect/SparseTensor/one_shot_bufferize_tensor_copy_insertion.mlir b/mlir/test/Dialect/SparseTensor/one_shot_bufferize_tensor_copy_insertion.mlir
index 8990fbf39b567a7..6c2292be161a531 100644
--- a/mlir/test/Dialect/SparseTensor/one_shot_bufferize_tensor_copy_insertion.mlir
+++ b/mlir/test/Dialect/SparseTensor/one_shot_bufferize_tensor_copy_insertion.mlir
@@ -15,7 +15,7 @@ func.func @bufferization_alloc_tensor() -> tensor<20x40xf32, #DCSR> {
return %1 : tensor<20x40xf32, #DCSR>
}
-!Filename = !llvm.ptr<i8>
+!Filename = !llvm.ptr
// CHECK-LABEL: func @sparse_tensor_new
// CHECK-FUNC-LABEL: func @sparse_tensor_new
func.func @sparse_tensor_new(%file: !Filename) -> tensor<20x40xf32, #DCSR> {
diff --git a/mlir/test/Dialect/SparseTensor/rewriting_for_codegen.mlir b/mlir/test/Dialect/SparseTensor/rewriting_for_codegen.mlir
index 93e802bc6065e42..579b0f20f5f6842 100644
--- a/mlir/test/Dialect/SparseTensor/rewriting_for_codegen.mlir
+++ b/mlir/test/Dialect/SparseTensor/rewriting_for_codegen.mlir
@@ -14,39 +14,39 @@
}>
// CHECK-LABEL: func.func @sparse_new(
-// CHECK-SAME: %[[A:.*]]: !llvm.ptr<i8>) -> tensor<?x?xf32, #sparse_tensor.encoding<{{{.*}}}>> {
-// CHECK: %[[COO:.*]] = sparse_tensor.new %[[A]] : !llvm.ptr<i8> to tensor<?x?xf32, #sparse_tensor.encoding<{{{.*}}}>>
+// CHECK-SAME: %[[A:.*]]: !llvm.ptr) -> tensor<?x?xf32, #sparse_tensor.encoding<{{{.*}}}>> {
+// CHECK: %[[COO:.*]] = sparse_tensor.new %[[A]] : !llvm.ptr to tensor<?x?xf32, #sparse_tensor.encoding<{{{.*}}}>>
// CHECK: %[[R:.*]] = sparse_tensor.convert %[[COO]]
// CHECK: bufferization.dealloc_tensor %[[COO]]
// CHECK: return %[[R]]
-func.func @sparse_new(%arg0: !llvm.ptr<i8>) -> tensor<?x?xf32, #CSR> {
- %0 = sparse_tensor.new %arg0 : !llvm.ptr<i8> to tensor<?x?xf32, #CSR>
+func.func @sparse_new(%arg0: !llvm.ptr) -> tensor<?x?xf32, #CSR> {
+ %0 = sparse_tensor.new %arg0 : !llvm.ptr to tensor<?x?xf32, #CSR>
return %0 : tensor<?x?xf32, #CSR>
}
// CHECK-LABEL: func.func @sparse_new_csc(
-// CHECK-SAME: %[[A:.*]]: !llvm.ptr<i8>) -> tensor<?x?xf32, #sparse_tensor.encoding<{{{.*}}}>> {
-// CHECK: %[[COO:.*]] = sparse_tensor.new %[[A]] : !llvm.ptr<i8> to tensor<?x?xf32, #sparse_tensor.encoding<{{{.*}}}>>
+// CHECK-SAME: %[[A:.*]]: !llvm.ptr) -> tensor<?x?xf32, #sparse_tensor.encoding<{{{.*}}}>> {
+// CHECK: %[[COO:.*]] = sparse_tensor.new %[[A]] : !llvm.ptr to tensor<?x?xf32, #sparse_tensor.encoding<{{{.*}}}>>
// CHECK: %[[R:.*]] = sparse_tensor.convert %[[COO]]
// CHECK: bufferization.dealloc_tensor %[[COO]]
// CHECK: return %[[R]]
-func.func @sparse_new_csc(%arg0: !llvm.ptr<i8>) -> tensor<?x?xf32, #CSC> {
- %0 = sparse_tensor.new %arg0 : !llvm.ptr<i8> to tensor<?x?xf32, #CSC>
+func.func @sparse_new_csc(%arg0: !llvm.ptr) -> tensor<?x?xf32, #CSC> {
+ %0 = sparse_tensor.new %arg0 : !llvm.ptr to tensor<?x?xf32, #CSC>
return %0 : tensor<?x?xf32, #CSC>
}
// CHECK-LABEL: func.func @sparse_new_coo(
-// CHECK-SAME: %[[A:.*]]: !llvm.ptr<i8>) -> tensor<?x?xf32, #sparse_tensor.encoding<{{{.*}}}>> {
-// CHECK: %[[COO:.*]] = sparse_tensor.new %[[A]] : !llvm.ptr<i8> to tensor<?x?xf32, #sparse_tensor.encoding<{{{.*}}}>>
+// CHECK-SAME: %[[A:.*]]: !llvm.ptr) -> tensor<?x?xf32, #sparse_tensor.encoding<{{{.*}}}>> {
+// CHECK: %[[COO:.*]] = sparse_tensor.new %[[A]] : !llvm.ptr to tensor<?x?xf32, #sparse_tensor.encoding<{{{.*}}}>>
// CHECK: return %[[COO]]
-func.func @sparse_new_coo(%arg0: !llvm.ptr<i8>) -> tensor<?x?xf32, #COO> {
- %0 = sparse_tensor.new %arg0 : !llvm.ptr<i8> to tensor<?x?xf32, #COO>
+func.func @sparse_new_coo(%arg0: !llvm.ptr) -> tensor<?x?xf32, #COO> {
+ %0 = sparse_tensor.new %arg0 : !llvm.ptr to tensor<?x?xf32, #COO>
return %0 : tensor<?x?xf32, #COO>
}
// CHECK-LABEL: func.func @sparse_out(
// CHECK-SAME: %[[A:.*]]: tensor<10x20xf32, #sparse_tensor.encoding<{{{.*}}}>>,
-// CHECK-SAME: %[[B:.*]]: !llvm.ptr<i8>) {
+// CHECK-SAME: %[[B:.*]]: !llvm.ptr) {
// CHECK-DAG: %[[C0:.*]] = arith.constant 0 : index
// CHECK-DAG: %[[C1:.*]] = arith.constant 1 : index
// CHECK-DAG: %[[C2:.*]] = arith.constant 2 : index
@@ -67,7 +67,7 @@ func.func @sparse_new_coo(%arg0: !llvm.ptr<i8>) -> tensor<?x?xf32, #COO> {
// CHECK: call @delSparseTensorWriter(%[[W]])
// CHECK: return
// CHECK: }
-func.func @sparse_out( %arg0: tensor<10x20xf32, #CSR>, %arg1: !llvm.ptr<i8>) -> () {
- sparse_tensor.out %arg0, %arg1 : tensor<10x20xf32, #CSR>, !llvm.ptr<i8>
+func.func @sparse_out( %arg0: tensor<10x20xf32, #CSR>, %arg1: !llvm.ptr) -> () {
+ sparse_tensor.out %arg0, %arg1 : tensor<10x20xf32, #CSR>, !llvm.ptr
return
}
diff --git a/mlir/test/Dialect/SparseTensor/roundtrip.mlir b/mlir/test/Dialect/SparseTensor/roundtrip.mlir
index 17ae8c065945a1b..5b733ecbe7874db 100644
--- a/mlir/test/Dialect/SparseTensor/roundtrip.mlir
+++ b/mlir/test/Dialect/SparseTensor/roundtrip.mlir
@@ -3,11 +3,11 @@
#SparseVector = #sparse_tensor.encoding<{map = (d0) -> (d0 : compressed)}>
// CHECK-LABEL: func @sparse_new(
-// CHECK-SAME: %[[A:.*]]: !llvm.ptr<i8>)
-// CHECK: %[[T:.*]] = sparse_tensor.new %[[A]] : !llvm.ptr<i8> to tensor<128xf64, #{{.*}}>
+// CHECK-SAME: %[[A:.*]]: !llvm.ptr)
+// CHECK: %[[T:.*]] = sparse_tensor.new %[[A]] : !llvm.ptr to tensor<128xf64, #{{.*}}>
// CHECK: return %[[T]] : tensor<128xf64, #{{.*}}>
-func.func @sparse_new(%arg0: !llvm.ptr<i8>) -> tensor<128xf64, #SparseVector> {
- %0 = sparse_tensor.new %arg0 : !llvm.ptr<i8> to tensor<128xf64, #SparseVector>
+func.func @sparse_new(%arg0: !llvm.ptr) -> tensor<128xf64, #SparseVector> {
+ %0 = sparse_tensor.new %arg0 : !llvm.ptr to tensor<128xf64, #SparseVector>
return %0 : tensor<128xf64, #SparseVector>
}
@@ -401,11 +401,11 @@ func.func @sparse_compression(%values: memref<?xf64>,
// CHECK-LABEL: func @sparse_out(
// CHECK-SAME: %[[A:.*]]: tensor<?x?xf64, #sparse_tensor.encoding<{{.*}}>>,
-// CHECK-SAME: %[[B:.*]]: !llvm.ptr<i8>)
-// CHECK: sparse_tensor.out %[[A]], %[[B]] : tensor<?x?xf64, #sparse_tensor.encoding<{{.*}}>>, !llvm.ptr<i8>
+// CHECK-SAME: %[[B:.*]]: !llvm.ptr)
+// CHECK: sparse_tensor.out %[[A]], %[[B]] : tensor<?x?xf64, #sparse_tensor.encoding<{{.*}}>>, !llvm.ptr
// CHECK: return
-func.func @sparse_out(%arg0: tensor<?x?xf64, #SparseMatrix>, %arg1: !llvm.ptr<i8>) {
- sparse_tensor.out %arg0, %arg1 : tensor<?x?xf64, #SparseMatrix>, !llvm.ptr<i8>
+func.func @sparse_out(%arg0: tensor<?x?xf64, #SparseMatrix>, %arg1: !llvm.ptr) {
+ sparse_tensor.out %arg0, %arg1 : tensor<?x?xf64, #SparseMatrix>, !llvm.ptr
return
}
diff --git a/mlir/test/mlir-cpu-runner/simple.mlir b/mlir/test/mlir-cpu-runner/simple.mlir
index 2734e499b6cab35..38d9dcaf553714f 100644
--- a/mlir/test/mlir-cpu-runner/simple.mlir
+++ b/mlir/test/mlir-cpu-runner/simple.mlir
@@ -15,8 +15,8 @@
// Declarations of C library functions.
llvm.func @logbf(f32) -> f32
-llvm.func @malloc(i64) -> !llvm.ptr<i8>
-llvm.func @free(!llvm.ptr<i8>)
+llvm.func @malloc(i64) -> !llvm.ptr
+llvm.func @free(!llvm.ptr)
// Check that a simple function with a nested call works.
llvm.func @main() -> f32 {
@@ -27,29 +27,27 @@ llvm.func @main() -> f32 {
// CHECK: 8.000000e+00
// Helper typed functions wrapping calls to "malloc" and "free".
-llvm.func @allocation() -> !llvm.ptr<f32> {
+llvm.func @allocation() -> !llvm.ptr {
%0 = llvm.mlir.constant(4 : index) : i64
- %1 = llvm.call @malloc(%0) : (i64) -> !llvm.ptr<i8>
- %2 = llvm.bitcast %1 : !llvm.ptr<i8> to !llvm.ptr<f32>
- llvm.return %2 : !llvm.ptr<f32>
+ %1 = llvm.call @malloc(%0) : (i64) -> !llvm.ptr
+ llvm.return %1 : !llvm.ptr
}
-llvm.func @deallocation(%arg0: !llvm.ptr<f32>) {
- %0 = llvm.bitcast %arg0 : !llvm.ptr<f32> to !llvm.ptr<i8>
- llvm.call @free(%0) : (!llvm.ptr<i8>) -> ()
+llvm.func @deallocation(%arg0: !llvm.ptr) {
+ llvm.call @free(%arg0) : (!llvm.ptr) -> ()
llvm.return
}
// Check that allocation and deallocation works, and that a custom entry point
// works.
llvm.func @foo() -> f32 {
- %0 = llvm.call @allocation() : () -> !llvm.ptr<f32>
+ %0 = llvm.call @allocation() : () -> !llvm.ptr
%1 = llvm.mlir.constant(0 : index) : i64
%2 = llvm.mlir.constant(1.234000e+03 : f32) : f32
- %3 = llvm.getelementptr %0[%1] : (!llvm.ptr<f32>, i64) -> !llvm.ptr<f32>
- llvm.store %2, %3 : !llvm.ptr<f32>
- %4 = llvm.getelementptr %0[%1] : (!llvm.ptr<f32>, i64) -> !llvm.ptr<f32>
- %5 = llvm.load %4 : !llvm.ptr<f32>
- llvm.call @deallocation(%0) : (!llvm.ptr<f32>) -> ()
+ %3 = llvm.getelementptr %0[%1] : (!llvm.ptr, i64) -> !llvm.ptr, f32
+ llvm.store %2, %3 : f32, !llvm.ptr
+ %4 = llvm.getelementptr %0[%1] : (!llvm.ptr, i64) -> !llvm.ptr, f32
+ %5 = llvm.load %4 : !llvm.ptr -> f32
+ llvm.call @deallocation(%0) : (!llvm.ptr) -> ()
llvm.return %5 : f32
}
// NOMAIN: 1.234000e+03
diff --git a/mlir/test/mlir-cpu-runner/x86-varargs.mlir b/mlir/test/mlir-cpu-runner/x86-varargs.mlir
index f3f4322ce87975e..de1b723f461e52a 100644
--- a/mlir/test/mlir-cpu-runner/x86-varargs.mlir
+++ b/mlir/test/mlir-cpu-runner/x86-varargs.mlir
@@ -36,31 +36,29 @@ llvm.func @foo(%arg0: i32, ...) -> i32 {
%10 = llvm.mlir.constant(0 : i64) : i64
%11 = llvm.mlir.constant(0 : i64) : i64
%12 = llvm.mlir.constant(1 : i32) : i32
- %13 = llvm.alloca %12 x !llvm.array<1 x struct<"struct.va_list", (i32, i32, ptr<i8>, ptr<i8>)>> {alignment = 8 : i64} : (i32) -> !llvm.ptr<array<1 x struct<"struct.va_list", (i32, i32, ptr<i8>, ptr<i8>)>>>
- %14 = llvm.bitcast %13 : !llvm.ptr<array<1 x struct<"struct.va_list", (i32, i32, ptr<i8>, ptr<i8>)>>> to !llvm.ptr<i8>
- llvm.intr.vastart %14 : !llvm.ptr<i8>
- %15 = llvm.getelementptr %13[%11, %10, 0] : (!llvm.ptr<array<1 x struct<"struct.va_list", (i32, i32, ptr<i8>, ptr<i8>)>>>, i64, i64) -> !llvm.ptr<i32>
- %16 = llvm.load %15 : !llvm.ptr<i32>
+ %13 = llvm.alloca %12 x !llvm.array<1 x struct<"struct.va_list", (i32, i32, ptr, ptr)>> {alignment = 8 : i64} : (i32) -> !llvm.ptr
+ llvm.intr.vastart %13 : !llvm.ptr
+ %15 = llvm.getelementptr %13[%11, %10, 0] : (!llvm.ptr, i64, i64) -> !llvm.ptr, !llvm.array<1 x struct<"struct.va_list", (i32, i32, ptr, ptr)>>
+ %16 = llvm.load %15 : !llvm.ptr -> i32
%17 = llvm.icmp "ult" %16, %8 : i32
llvm.cond_br %17, ^bb1, ^bb2
^bb1: // pred: ^bb0
- %18 = llvm.getelementptr %13[%7, %6, 3] : (!llvm.ptr<array<1 x struct<"struct.va_list", (i32, i32, ptr<i8>, ptr<i8>)>>>, i64, i64) -> !llvm.ptr<ptr<i8>>
- %19 = llvm.load %18 : !llvm.ptr<ptr<i8>>
+ %18 = llvm.getelementptr %13[%7, %6, 3] : (!llvm.ptr, i64, i64) -> !llvm.ptr, !llvm.array<1 x struct<"struct.va_list", (i32, i32, ptr, ptr)>>
+ %19 = llvm.load %18 : !llvm.ptr -> !llvm.ptr
%20 = llvm.zext %16 : i32 to i64
- %21 = llvm.getelementptr %19[%20] : (!llvm.ptr<i8>, i64) -> !llvm.ptr<i8>
+ %21 = llvm.getelementptr %19[%20] : (!llvm.ptr, i64) -> !llvm.ptr, i8
%22 = llvm.add %16, %4 : i32
- llvm.store %22, %15 : !llvm.ptr<i32>
- llvm.br ^bb3(%21 : !llvm.ptr<i8>)
+ llvm.store %22, %15 : i32, !llvm.ptr
+ llvm.br ^bb3(%21 : !llvm.ptr)
^bb2: // pred: ^bb0
- %23 = llvm.getelementptr %13[%3, %2, 2] : (!llvm.ptr<array<1 x struct<"struct.va_list", (i32, i32, ptr<i8>, ptr<i8>)>>>, i64, i64) -> !llvm.ptr<ptr<i8>>
- %24 = llvm.load %23 : !llvm.ptr<ptr<i8>>
- %25 = llvm.getelementptr %24[%0] : (!llvm.ptr<i8>, i64) -> !llvm.ptr<i8>
- llvm.store %25, %23 : !llvm.ptr<ptr<i8>>
- llvm.br ^bb3(%24 : !llvm.ptr<i8>)
-^bb3(%26: !llvm.ptr<i8>): // 2 preds: ^bb1, ^bb2
- %27 = llvm.bitcast %26 : !llvm.ptr<i8> to !llvm.ptr<i32>
- %28 = llvm.load %27 : !llvm.ptr<i32>
- llvm.intr.vaend %14 : !llvm.ptr<i8>
+ %23 = llvm.getelementptr %13[%3, %2, 2] : (!llvm.ptr, i64, i64) -> !llvm.ptr, !llvm.array<1 x struct<"struct.va_list", (i32, i32, ptr, ptr)>>
+ %24 = llvm.load %23 : !llvm.ptr -> !llvm.ptr
+ %25 = llvm.getelementptr %24[%0] : (!llvm.ptr, i64) -> !llvm.ptr, i8
+ llvm.store %25, %23 : !llvm.ptr, !llvm.ptr
+ llvm.br ^bb3(%24 : !llvm.ptr)
+^bb3(%26: !llvm.ptr): // 2 preds: ^bb1, ^bb2
+ %28 = llvm.load %26 : !llvm.ptr -> i32
+ llvm.intr.vaend %13 : !llvm.ptr
llvm.return %28 : i32
}
diff --git a/mlir/test/python/dialects/gpu/module-to-binary-nvvm.py b/mlir/test/python/dialects/gpu/module-to-binary-nvvm.py
index 70c08ceb7a6f2d3..1c2eb652e71f910 100644
--- a/mlir/test/python/dialects/gpu/module-to-binary-nvvm.py
+++ b/mlir/test/python/dialects/gpu/module-to-binary-nvvm.py
@@ -22,8 +22,8 @@ def testGPUToLLVMBin():
r"""
module attributes {gpu.container_module} {
gpu.module @kernel_module1 [#nvvm.target<chip = "sm_70">] {
- llvm.func @kernel(%arg0: i32, %arg1: !llvm.ptr<f32>,
- %arg2: !llvm.ptr<f32>, %arg3: i64, %arg4: i64,
+ llvm.func @kernel(%arg0: i32, %arg1: !llvm.ptr,
+ %arg2: !llvm.ptr, %arg3: i64, %arg4: i64,
%arg5: i64) attributes {gpu.kernel} {
llvm.return
}
@@ -47,8 +47,8 @@ def testGPUToASMBin():
r"""
module attributes {gpu.container_module} {
gpu.module @kernel_module2 [#nvvm.target<flags = {fast}>, #nvvm.target] {
- llvm.func @kernel(%arg0: i32, %arg1: !llvm.ptr<f32>,
- %arg2: !llvm.ptr<f32>, %arg3: i64, %arg4: i64,
+ llvm.func @kernel(%arg0: i32, %arg1: !llvm.ptr,
+ %arg2: !llvm.ptr, %arg3: i64, %arg4: i64,
%arg5: i64) attributes {gpu.kernel} {
llvm.return
}
diff --git a/mlir/test/python/dialects/gpu/module-to-binary-rocdl.py b/mlir/test/python/dialects/gpu/module-to-binary-rocdl.py
index fad088cbd6d893b..f5ee2856f39502d 100644
--- a/mlir/test/python/dialects/gpu/module-to-binary-rocdl.py
+++ b/mlir/test/python/dialects/gpu/module-to-binary-rocdl.py
@@ -22,8 +22,8 @@ def testGPUToLLVMBin():
r"""
module attributes {gpu.container_module} {
gpu.module @kernel_module1 [#rocdl.target<chip = "gfx90a">] {
- llvm.func @kernel(%arg0: i32, %arg1: !llvm.ptr<f32>,
- %arg2: !llvm.ptr<f32>, %arg3: i64, %arg4: i64,
+ llvm.func @kernel(%arg0: i32, %arg1: !llvm.ptr,
+ %arg2: !llvm.ptr, %arg3: i64, %arg4: i64,
%arg5: i64) attributes {gpu.kernel} {
llvm.return
}
@@ -47,8 +47,8 @@ def testGPUToASMBin():
r"""
module attributes {gpu.container_module} {
gpu.module @kernel_module2 [#rocdl.target<flags = {fast}>, #rocdl.target] {
- llvm.func @kernel(%arg0: i32, %arg1: !llvm.ptr<f32>,
- %arg2: !llvm.ptr<f32>, %arg3: i64, %arg4: i64,
+ llvm.func @kernel(%arg0: i32, %arg1: !llvm.ptr,
+ %arg2: !llvm.ptr, %arg3: i64, %arg4: i64,
%arg5: i64) attributes {gpu.kernel} {
llvm.return
}
More information about the Mlir-commits
mailing list