The target clause in OpenMP is used to offload loop kernels to accelarator
peripeherals. target's 'map' clause is used to move data from and to the
accelarator. When the data is SVE type, it may not be suitable because of
various reasons i.e. the two SVE targets may not agree on vector size or
some targets don't support variable vector size. This makes SVE unsuitable
for use in OMP's 'map' clause. This patch diagnoses all such cases and issues
an error where SVE types are not suitable.
Co-authored-by: Andrea Corallo <andrea.cora...@arm.com>
gcc/ChangeLog:
* target.h (type_context_kind): Add new context kinds for target
clauses.
* config/aarch64/aarch64-sve-builtins.cc (verify_type_context): Diagnose
SVE types for a given OpenMP context.
* gimplify.cc (omp_notice_variable): Diagnose implicitly-mapped SVE
objects in OpenMP regions.
(gimplify_scan_omp_clauses): Diagnose SVE types for various target
clauses.
gcc/testsuite/ChangeLog:
* gcc.target/aarch64/sve/omp/offload-1.c: New test.
* gcc.target/aarch64/sve/omp/offload-2.c: Likewise.
* gcc.target/aarch64/sve/omp/offload-parallel-loop.c: Likewise.
* gcc.target/aarch64/sve/omp/offload-parallel.c: Likewise.
* gcc.target/aarch64/sve/omp/offload-simd.c: Likewise.
* gcc.target/aarch64/sve/omp/offload-teams-distribute-simd.c: Likewise.
* gcc.target/aarch64/sve/omp/offload-teams-distribute.c: Likewise.
* gcc.target/aarch64/sve/omp/offload-teams-loop.c: Likewise.
* gcc.target/aarch64/sve/omp/offload-teams.c: Likewise.
* gcc.target/aarch64/sve/omp/target-device.c: Likewise.
* gcc.target/aarch64/sve/omp/target-link.c: Likewise.
---
gcc/config/aarch64/aarch64-sve-builtins.cc | 31 +++
gcc/gimplify.cc | 34 ++-
gcc/target.h | 19 +-
.../gcc.target/aarch64/sve/omp/offload-1.c | 237 ++++++++++++++++++
.../gcc.target/aarch64/sve/omp/offload-2.c | 198 +++++++++++++++
.../aarch64/sve/omp/offload-parallel-loop.c | 236 +++++++++++++++++
.../aarch64/sve/omp/offload-parallel.c | 195 ++++++++++++++
.../gcc.target/aarch64/sve/omp/offload-simd.c | 236 +++++++++++++++++
.../sve/omp/offload-teams-distribute-simd.c | 237 ++++++++++++++++++
.../sve/omp/offload-teams-distribute.c | 236 +++++++++++++++++
.../aarch64/sve/omp/offload-teams-loop.c | 237 ++++++++++++++++++
.../aarch64/sve/omp/offload-teams.c | 195 ++++++++++++++
.../aarch64/sve/omp/target-device.c | 97 +++++++
.../gcc.target/aarch64/sve/omp/target-link.c | 48 ++++
14 files changed, 2234 insertions(+), 2 deletions(-)
create mode 100644 gcc/testsuite/gcc.target/aarch64/sve/omp/offload-1.c
create mode 100644 gcc/testsuite/gcc.target/aarch64/sve/omp/offload-2.c
create mode 100644
gcc/testsuite/gcc.target/aarch64/sve/omp/offload-parallel-loop.c
create mode 100644 gcc/testsuite/gcc.target/aarch64/sve/omp/offload-parallel.c
create mode 100644 gcc/testsuite/gcc.target/aarch64/sve/omp/offload-simd.c
create mode 100644
gcc/testsuite/gcc.target/aarch64/sve/omp/offload-teams-distribute-simd.c
create mode 100644
gcc/testsuite/gcc.target/aarch64/sve/omp/offload-teams-distribute.c
create mode 100644
gcc/testsuite/gcc.target/aarch64/sve/omp/offload-teams-loop.c
create mode 100644 gcc/testsuite/gcc.target/aarch64/sve/omp/offload-teams.c
create mode 100644 gcc/testsuite/gcc.target/aarch64/sve/omp/target-device.c
create mode 100644 gcc/testsuite/gcc.target/aarch64/sve/omp/target-link.c
diff --git a/gcc/config/aarch64/aarch64-sve-builtins.cc
b/gcc/config/aarch64/aarch64-sve-builtins.cc
index f3983a123e3..ee1064c3bb7 100644
--- a/gcc/config/aarch64/aarch64-sve-builtins.cc
+++ b/gcc/config/aarch64/aarch64-sve-builtins.cc
@@ -5000,6 +5000,29 @@ bool
verify_type_context (location_t loc, type_context_kind context,
const_tree type, bool silent_p)
{
+ if (aarch64_sve::builtin_type_p (type)
+ || (POINTER_TYPE_P (type)
+ && aarch64_sve::builtin_type_p (TREE_TYPE (type))))
+ switch (context)
+ {
+ case TCTX_OMP_MAP:
+ error_at (loc, "SVE type %qT not allowed in map clause", type);
+ return false;
+ case TCTX_OMP_MAP_IMP_REF:
+ return false;
+ case TCTX_OMP_PRIVATE:
+ error_at (loc, "SVE type %qT not allowed in target private clause",
type);
+ return false;
+ case TCTX_OMP_FIRSTPRIVATE:
+ error_at (loc, "SVE type %qT not allowed in target firstprivate
clause", type);
+ return false;
+ case TCTX_OMP_DEVICE_ADDR:
+ error_at (loc, "SVE type %qT not allowed in target device clauses",
type);
+ return false;
+ default:
+ break;
+ }
+
if (!sizeless_type_p (type))
return true;
@@ -5060,6 +5083,14 @@ verify_type_context (location_t loc, type_context_kind
context,
if (!silent_p)
error_at (loc, "capture by copy of SVE type %qT", type);
return false;
+
+ case TCTX_OMP_MAP:
+ case TCTX_OMP_MAP_IMP_REF:
+ case TCTX_OMP_PRIVATE:
+ case TCTX_OMP_FIRSTPRIVATE:
+ case TCTX_OMP_DEVICE_ADDR:
+ default:
+ break;
}
gcc_unreachable ();
}
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index d87eb433395..dc958d2f55d 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -8349,11 +8349,13 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree
decl, bool in_code)
| GOVD_MAP_ALLOC_ONLY)) == flags)
{
tree type = TREE_TYPE (decl);
+ location_t dummy = UNKNOWN_LOCATION;
if (gimplify_omp_ctxp->target_firstprivatize_array_bases
&& omp_privatize_by_reference (decl))
type = TREE_TYPE (type);
- if (!omp_mappable_type (type))
+ if (!omp_mappable_type (type)
+ || !verify_type_context (dummy, TCTX_OMP_MAP_IMP_REF, type))
{
error ("%qD referenced in target region does not have "
"a mappable type", decl);
@@ -12083,6 +12085,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
unsigned int flags;
tree decl;
auto_vec<omp_addr_token *, 10> addr_tokens;
+ tree op = NULL_TREE;
+ location_t loc = OMP_CLAUSE_LOCATION (c);
if (grp_end && c == OMP_CLAUSE_CHAIN (grp_end))
{
@@ -12090,6 +12094,34 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
grp_end = NULL_TREE;
}
+ if (code == OMP_TARGET || code == OMP_TARGET_DATA
+ || code == OMP_TARGET_ENTER_DATA || code == OMP_TARGET_EXIT_DATA)
+ /* Do some target-specific type checks for map operands. */
+ switch (OMP_CLAUSE_CODE (c))
+ {
+ case OMP_CLAUSE_MAP:
+ op = OMP_CLAUSE_OPERAND (c, 0);
+ verify_type_context (loc, TCTX_OMP_MAP, TREE_TYPE (op));
+ break;
+ case OMP_CLAUSE_PRIVATE:
+ op = OMP_CLAUSE_OPERAND (c, 0);
+ verify_type_context (loc, TCTX_OMP_PRIVATE, TREE_TYPE (op));
+ break;
+ case OMP_CLAUSE_FIRSTPRIVATE:
+ op = OMP_CLAUSE_OPERAND (c, 0);
+ verify_type_context (loc, TCTX_OMP_FIRSTPRIVATE, TREE_TYPE (op));
+ break;
+ case OMP_CLAUSE_IS_DEVICE_PTR:
+ case OMP_CLAUSE_USE_DEVICE_ADDR:
+ case OMP_CLAUSE_USE_DEVICE_PTR:
+ case OMP_CLAUSE_HAS_DEVICE_ADDR:
+ op = OMP_CLAUSE_OPERAND (c, 0);
+ verify_type_context (loc, TCTX_OMP_DEVICE_ADDR, TREE_TYPE (op));
+ break;
+ default:
+ break;
+ }
+
switch (OMP_CLAUSE_CODE (c))
{
case OMP_CLAUSE_PRIVATE:
diff --git a/gcc/target.h b/gcc/target.h
index c1f99b97b86..9cebd354fdb 100644
--- a/gcc/target.h
+++ b/gcc/target.h
@@ -271,7 +271,24 @@ enum type_context_kind {
TCTX_EXCEPTIONS,
/* Capturing objects of type T by value in a closure. */
- TCTX_CAPTURE_BY_COPY
+ TCTX_CAPTURE_BY_COPY,
+
+ /* Objects of type T appearing in OpenMP map clause. */
+ TCTX_OMP_MAP,
+
+ /* Objects of type T appearing in OpenMP target region
+ without explicit map. */
+ TCTX_OMP_MAP_IMP_REF,
+
+ /* Objects of type T appearing in OpenMP private clause. */
+ TCTX_OMP_PRIVATE,
+
+ /* Objects of type T appearing in OpenMP firstprivate clause. */
+ TCTX_OMP_FIRSTPRIVATE,
+
+ /* Objects of type T appearing in OpenMP device clauses. */
+ TCTX_OMP_DEVICE_ADDR
+
};
enum poly_value_estimate_kind
diff --git a/gcc/testsuite/gcc.target/aarch64/sve/omp/offload-1.c
b/gcc/testsuite/gcc.target/aarch64/sve/omp/offload-1.c
new file mode 100644
index 00000000000..20dd478e079
--- /dev/null
+++ b/gcc/testsuite/gcc.target/aarch64/sve/omp/offload-1.c
@@ -0,0 +1,237 @@
+/* { dg-do compile } */
+/* { dg-options "-msve-vector-bits=256 -std=gnu99 -fopenmp -O2
-fdump-tree-ompexp" } */
+
+#include <arm_sve.h>
+
+#define N 256
+
+#ifndef CONSTRUCT
+#define CONSTRUCT
+#endif
+
+svint32_t
+__attribute__ ((noinline))
+omp_target ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target CONSTRUCT
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc); /* { dg-error {'va'
referenced in target region does not have a mappable type} } */
+ }
+
+ return va;
+}
+
+svint32_t
+__attribute__ ((noinline))
+omp_target_data_map_1 ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target CONSTRUCT map(to: b, c) map(from: va) /* { dg-error {SVE
type 'svint32_t' not allowed in map clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc);
+ }
+
+ return va;
+}
+
+svint32_t
+__attribute__ ((noinline))
+omp_target_data_map_2 ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target CONSTRUCT map(to: b, c) map(from: va) /* { dg-error {SVE
type 'svint32_t' not allowed in map clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc);
+ }
+
+#pragma omp target CONSTRUCT map(to: b, c) map(tofrom: va) /* { dg-error {SVE
type 'svint32_t' not allowed in map clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, va);
+ va = svadd_s32_z (svptrue_b32 (), vc, va);
+ }
+
+ return va;
+}
+
+svint32_t
+__attribute__ ((noinline))
+omp_target_map_data_enter_exit ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target enter data map(to: b, c)
+
+#pragma omp target CONSTRUCT map(from: va) /* { dg-error {SVE type 'svint32_t'
not allowed in map clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc);
+ }
+
+#pragma omp target CONSTRUCT map(to: va) /* { dg-error {SVE type 'svint32_t'
not allowed in map clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, va);
+ va = svadd_s32_z (svptrue_b32 (), vc, va);
+ }
+
+#pragma omp target exit data map(from: va) /* { dg-error {SVE type 'svint32_t'
not allowed in map clause} } */
+
+ return va;
+}
+
+svint32_t
+__attribute__ ((noinline))
+omp_target_map_data_alloc_update ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target data map(to: b, c) map(alloc: va) /* { dg-error {SVE type
'svint32_t' not allowed in map clause} } */
+{
+#pragma omp target CONSTRUCT
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc); /* { dg-error {'va'
referenced in target region does not have a mappable type} } */
+ }
+
+/* Update va on the host from target. */
+#pragma omp target update from(va)
+
+#pragma omp target CONSTRUCT map(from: va) /* { dg-error {SVE type 'svint32_t'
not allowed in map clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, va);
+ va = svadd_s32_z (svptrue_b32 (), vc, va);
+ }
+}
+ return va;
+}
+
+int64_t __attribute__ ((noinline))
+omp_target_private ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int64_t res;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target CONSTRUCT private (va, vb, vc) map (to: b, c) map (from:
res) /* { dg-error {SVE type 'svint32_t' not allowed in target private clause}
} */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b);
+ vc = svld1_s32 (svptrue_b32 (), c);
+ va = svadd_s32_z (svptrue_b32 (), vb, vc);
+ res = svaddv_s32 (svptrue_b32 (), va);
+ }
+
+ return res;
+}
+
+int64_t __attribute__ ((noinline))
+omp_target_firstprivate (svbool_t vp)
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int64_t res;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target CONSTRUCT firstprivate (vp) map (to: b, c) map (from:
res)/* { dg-error {SVE type 'svbool_t' not allowed in target firstprivate
clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (vp, b); /* { dg-error {'vb' referenced in target region
does not have a mappable type} } */
+ vc = svld1_s32 (vp, c); /* { dg-error {'vc' referenced in target region
does not have a mappable type} } */
+ va = svadd_s32_z (vp, vb, vc); /* { dg-error {'va' referenced in target
region does not have a mappable type} } */
+ res = svaddv_s32 (svptrue_b32 (), va);
+ }
+
+ return res;
+}
diff --git a/gcc/testsuite/gcc.target/aarch64/sve/omp/offload-2.c
b/gcc/testsuite/gcc.target/aarch64/sve/omp/offload-2.c
new file mode 100644
index 00000000000..efb4d274de8
--- /dev/null
+++ b/gcc/testsuite/gcc.target/aarch64/sve/omp/offload-2.c
@@ -0,0 +1,198 @@
+/* { dg-do compile } */
+/* { dg-options "-msve-vector-bits=256 -std=gnu99 -fopenmp -O2
-fdump-tree-ompexp" } */
+
+#include <arm_sve.h>
+
+#define N 256
+
+#ifndef CONSTRUCT
+#define CONSTRUCT
+#endif
+
+svint32_t
+__attribute__ ((noinline))
+omp_target ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target CONSTRUCT
+ {
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc); /* { dg-error {'va'
referenced in target region does not have a mappable type} } */
+ }
+ }
+
+ return va;
+}
+
+svint32_t
+__attribute__ ((noinline))
+omp_target_data_map_1 ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target CONSTRUCT map(to: b, c) map(from: va) /* { dg-error {SVE
type 'svint32_t' not allowed in map clause} } */
+ {
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc);
+ }
+ }
+
+ return va;
+}
+
+svint32_t
+__attribute__ ((noinline))
+omp_target_data_map_2 ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+
+#pragma omp target CONSTRUCT map(to: b, c) map(from: va) /* { dg-error {SVE
type 'svint32_t' not allowed in map clause} } */
+ {
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc);
+ }
+ }
+
+#pragma omp target CONSTRUCT map(to: b, c) map(tofrom: va) /* { dg-error {SVE
type 'svint32_t' not allowed in map clause} } */
+ {
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, va);
+ va = svadd_s32_z (svptrue_b32 (), vc, va);
+ }
+ }
+
+ return va;
+}
+
+svint32_t
+__attribute__ ((noinline))
+omp_target_map_data_enter_exit ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target enter data map(to: b, c)
+
+#pragma omp target CONSTRUCT map(from: va) /* { dg-error {SVE type 'svint32_t'
not allowed in map clause} } */
+ {
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc);
+ }
+ }
+
+#pragma omp target CONSTRUCT map(to: va) /* { dg-error {SVE type 'svint32_t'
not allowed in map clause} } */
+ {
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, va);
+ va = svadd_s32_z (svptrue_b32 (), vc, va);
+ }
+ }
+
+#pragma omp target exit data map(from: va) /* { dg-error {SVE type 'svint32_t'
not allowed in map clause} } */
+
+ return va;
+}
+
+svint32_t
+__attribute__ ((noinline))
+omp_target_map_data_alloc_update ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target data map(to: b, c) map(alloc: va) /* { dg-error {SVE type
'svint32_t' not allowed in map clause} } */
+{
+#pragma omp target CONSTRUCT
+ {
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc); /* { dg-error {'va'
referenced in target region does not have a mappable type} } */
+ }
+ }
+
+/* Update va on the host from target. */
+#pragma omp target update from(va)
+
+#pragma omp target CONSTRUCT map(from: va) /* { dg-error {SVE type 'svint32_t'
not allowed in map clause} } */
+ {
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, va);
+ va = svadd_s32_z (svptrue_b32 (), vc, va);
+ }
+ }
+}
+ return va;
+}
diff --git a/gcc/testsuite/gcc.target/aarch64/sve/omp/offload-parallel-loop.c
b/gcc/testsuite/gcc.target/aarch64/sve/omp/offload-parallel-loop.c
new file mode 100644
index 00000000000..4c6a0d4d96a
--- /dev/null
+++ b/gcc/testsuite/gcc.target/aarch64/sve/omp/offload-parallel-loop.c
@@ -0,0 +1,236 @@
+/* { dg-do compile } */
+/* { dg-options "-msve-vector-bits=256 -std=gnu99 -fopenmp -O2
-fdump-tree-ompexp" } */
+
+#include <arm_sve.h>
+
+#define N 256
+#define CONSTRUCT parallel loop
+
+svint32_t
+__attribute__ ((noinline))
+omp_target ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target CONSTRUCT
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc); /* { dg-error {'va'
referenced in target region does not have a mappable type} } */
+ }
+
+ return va;
+}
+
+svint32_t
+__attribute__ ((noinline))
+omp_target_data_map_1 ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target CONSTRUCT map(to: b, c) map(from: va) /* { dg-error {SVE
type 'svint32_t' not allowed in map clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc);
+ }
+
+ return va;
+}
+
+svint32_t
+__attribute__ ((noinline))
+omp_target_data_map_2 ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target CONSTRUCT map(to: b, c) map(from: va) /* { dg-error {SVE
type 'svint32_t' not allowed in map clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc);
+ }
+
+#pragma omp target CONSTRUCT map(to: b, c) map(tofrom: va) /* { dg-error {SVE
type 'svint32_t' not allowed in map clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, va);
+ va = svadd_s32_z (svptrue_b32 (), vc, va);
+ }
+
+ return va;
+}
+
+svint32_t
+__attribute__ ((noinline))
+omp_target_map_data_enter_exit ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target enter data map(to: b, c)
+
+#pragma omp target CONSTRUCT map(from: va) /* { dg-error {SVE type 'svint32_t'
not allowed in map clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc);
+ }
+
+#pragma omp target CONSTRUCT map(to: va) /* { dg-error {SVE type 'svint32_t'
not allowed in map clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, va);
+ va = svadd_s32_z (svptrue_b32 (), vc, va);
+ }
+
+#pragma omp target exit data map(from: va) /* { dg-error {SVE type 'svint32_t'
not allowed in map clause} } */
+
+ return va;
+}
+
+svint32_t
+__attribute__ ((noinline))
+omp_target_map_data_alloc_update ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target data map(to: b, c) map(alloc: va) /* { dg-error {SVE type
'svint32_t' not allowed in map clause} } */
+{
+#pragma omp target CONSTRUCT
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc); /* { dg-error {'va'
referenced in target region does not have a mappable type} } */
+ }
+
+/* Update va on the host from target. */
+#pragma omp target update from(va)
+
+#pragma omp target CONSTRUCT map(from: va) /* { dg-error {SVE type 'svint32_t'
not allowed in map clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, va);
+ va = svadd_s32_z (svptrue_b32 (), vc, va);
+ }
+}
+ return va;
+}
+
+int64_t __attribute__ ((noinline))
+omp_target_private ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int64_t res;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+/* Combined construct scenario: here private applies to the parallel loop
+ construct, so no error. */
+#pragma omp target CONSTRUCT private (va, vb, vc) map (to: b, c) map (from:
res)
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b);
+ vc = svld1_s32 (svptrue_b32 (), c);
+ va = svadd_s32_z (svptrue_b32 (), vb, vc);
+ res = svaddv_s32 (svptrue_b32 (), va);
+ }
+
+ return res;
+}
+
+int64_t __attribute__ ((noinline))
+omp_target_firstprivate (svbool_t vp)
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int64_t res;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target CONSTRUCT firstprivate (vp) map (to: b, c) map (from:
res)/* { dg-error {SVE type 'svbool_t' not allowed in target firstprivate
clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (vp, b); /* { dg-error {'vb' referenced in target region
does not have a mappable type} } */
+ vc = svld1_s32 (vp, c); /* { dg-error {'vc' referenced in target region
does not have a mappable type} } */
+ va = svadd_s32_z (vp, vb, vc); /* { dg-error {'va' referenced in target
region does not have a mappable type} } */
+ res = svaddv_s32 (svptrue_b32 (), va);
+ }
+
+ return res;
+}
diff --git a/gcc/testsuite/gcc.target/aarch64/sve/omp/offload-parallel.c
b/gcc/testsuite/gcc.target/aarch64/sve/omp/offload-parallel.c
new file mode 100644
index 00000000000..39dcd39a5f5
--- /dev/null
+++ b/gcc/testsuite/gcc.target/aarch64/sve/omp/offload-parallel.c
@@ -0,0 +1,195 @@
+/* { dg-do compile } */
+/* { dg-options "-msve-vector-bits=256 -std=gnu99 -fopenmp -O2
-fdump-tree-ompexp" } */
+
+#include <arm_sve.h>
+
+#define CONSTRUCT parallel
+#define N 256
+
+svint32_t
+__attribute__ ((noinline))
+omp_target ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target CONSTRUCT
+ {
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc); /* { dg-error {'va'
referenced in target region does not have a mappable type} } */
+ }
+ }
+
+ return va;
+}
+
+svint32_t
+__attribute__ ((noinline))
+omp_target_data_map_1 ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target CONSTRUCT map(to: b, c) map(from: va) /* { dg-error {SVE
type 'svint32_t' not allowed in map clause} } */
+ {
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc);
+ }
+ }
+
+ return va;
+}
+
+svint32_t
+__attribute__ ((noinline))
+omp_target_data_map_2 ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+
+#pragma omp target CONSTRUCT map(to: b, c) map(from: va) /* { dg-error {SVE
type 'svint32_t' not allowed in map clause} } */
+ {
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc);
+ }
+ }
+
+#pragma omp target CONSTRUCT map(to: b, c) map(tofrom: va) /* { dg-error {SVE
type 'svint32_t' not allowed in map clause} } */
+ {
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, va);
+ va = svadd_s32_z (svptrue_b32 (), vc, va);
+ }
+ }
+
+ return va;
+}
+
+svint32_t
+__attribute__ ((noinline))
+omp_target_map_data_enter_exit ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target enter data map(to: b, c)
+
+#pragma omp target CONSTRUCT map(from: va) /* { dg-error {SVE type 'svint32_t'
not allowed in map clause} } */
+ {
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc);
+ }
+ }
+
+#pragma omp target CONSTRUCT map(to: va) /* { dg-error {SVE type 'svint32_t'
not allowed in map clause} } */
+ {
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, va);
+ va = svadd_s32_z (svptrue_b32 (), vc, va);
+ }
+ }
+
+#pragma omp target exit data map(from: va) /* { dg-error {SVE type 'svint32_t'
not allowed in map clause} } */
+
+ return va;
+}
+
+svint32_t
+__attribute__ ((noinline))
+omp_target_map_data_alloc_update ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target data map(to: b, c) map(alloc: va) /* { dg-error {SVE type
'svint32_t' not allowed in map clause} } */
+{
+#pragma omp target CONSTRUCT
+ {
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc); /* { dg-error {'va'
referenced in target region does not have a mappable type} } */
+ }
+ }
+
+/* Update va on the host from target. */
+#pragma omp target update from(va)
+
+#pragma omp target CONSTRUCT map(from: va) /* { dg-error {SVE type 'svint32_t'
not allowed in map clause} } */
+ {
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, va);
+ va = svadd_s32_z (svptrue_b32 (), vc, va);
+ }
+ }
+}
+ return va;
+}
diff --git a/gcc/testsuite/gcc.target/aarch64/sve/omp/offload-simd.c
b/gcc/testsuite/gcc.target/aarch64/sve/omp/offload-simd.c
new file mode 100644
index 00000000000..2bb2a884fcf
--- /dev/null
+++ b/gcc/testsuite/gcc.target/aarch64/sve/omp/offload-simd.c
@@ -0,0 +1,236 @@
+/* { dg-do compile } */
+/* { dg-options "-msve-vector-bits=256 -std=gnu99 -fopenmp -O2
-fdump-tree-ompexp" } */
+
+#include <arm_sve.h>
+
+#define N 256
+#define CONSTRUCT simd
+
+svint32_t
+__attribute__ ((noinline))
+omp_target ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target CONSTRUCT
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc); /* { dg-error {'va'
referenced in target region does not have a mappable type} } */
+ }
+
+ return va;
+}
+
+svint32_t
+__attribute__ ((noinline))
+omp_target_data_map_1 ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target CONSTRUCT map(to: b, c) map(from: va) /* { dg-error {SVE
type 'svint32_t' not allowed in map clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc);
+ }
+
+ return va;
+}
+
+svint32_t
+__attribute__ ((noinline))
+omp_target_data_map_2 ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target CONSTRUCT map(to: b, c) map(from: va) /* { dg-error {SVE
type 'svint32_t' not allowed in map clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc);
+ }
+
+#pragma omp target CONSTRUCT map(to: b, c) map(tofrom: va) /* { dg-error {SVE
type 'svint32_t' not allowed in map clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, va);
+ va = svadd_s32_z (svptrue_b32 (), vc, va);
+ }
+
+ return va;
+}
+
+svint32_t
+__attribute__ ((noinline))
+omp_target_map_data_enter_exit ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target enter data map(to: b, c)
+
+#pragma omp target CONSTRUCT map(from: va) /* { dg-error {SVE type 'svint32_t'
not allowed in map clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc);
+ }
+
+#pragma omp target CONSTRUCT map(to: va) /* { dg-error {SVE type 'svint32_t'
not allowed in map clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, va);
+ va = svadd_s32_z (svptrue_b32 (), vc, va);
+ }
+
+#pragma omp target exit data map(from: va) /* { dg-error {SVE type 'svint32_t'
not allowed in map clause} } */
+
+ return va;
+}
+
+svint32_t
+__attribute__ ((noinline))
+omp_target_map_data_alloc_update ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target data map(to: b, c) map(alloc: va) /* { dg-error {SVE type
'svint32_t' not allowed in map clause} } */
+{
+#pragma omp target CONSTRUCT
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc); /* { dg-error {'va'
referenced in target region does not have a mappable type} } */
+ }
+
+/* Update va on the host from target. */
+#pragma omp target update from(va)
+
+#pragma omp target CONSTRUCT map(from: va) /* { dg-error {SVE type 'svint32_t'
not allowed in map clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, va);
+ va = svadd_s32_z (svptrue_b32 (), vc, va);
+ }
+}
+ return va;
+}
+
+int64_t __attribute__ ((noinline))
+omp_target_private ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int64_t res;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+/* Combined construct scenario: here private applies to the simd construct so
+ no error. */
+#pragma omp target CONSTRUCT private (va, vb, vc) map (to: b, c) map (from:
res)
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b);
+ vc = svld1_s32 (svptrue_b32 (), c);
+ va = svadd_s32_z (svptrue_b32 (), vb, vc);
+ res = svaddv_s32 (svptrue_b32 (), va);
+ }
+
+ return res;
+}
+
+int64_t __attribute__ ((noinline))
+omp_target_firstprivate (svbool_t vp)
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int64_t res;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target CONSTRUCT firstprivate (vp) map (to: b, c) map (from:
res)/* { dg-error {SVE type 'svbool_t' not allowed in target firstprivate
clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (vp, b); /* { dg-error {'vb' referenced in target region
does not have a mappable type} } */
+ vc = svld1_s32 (vp, c); /* { dg-error {'vc' referenced in target region
does not have a mappable type} } */
+ va = svadd_s32_z (vp, vb, vc); /* { dg-error {'va' referenced in target
region does not have a mappable type} } */
+ res = svaddv_s32 (svptrue_b32 (), va);
+ }
+
+ return res;
+}
diff --git
a/gcc/testsuite/gcc.target/aarch64/sve/omp/offload-teams-distribute-simd.c
b/gcc/testsuite/gcc.target/aarch64/sve/omp/offload-teams-distribute-simd.c
new file mode 100644
index 00000000000..6a61883e80a
--- /dev/null
+++ b/gcc/testsuite/gcc.target/aarch64/sve/omp/offload-teams-distribute-simd.c
@@ -0,0 +1,237 @@
+/* { dg-do compile } */
+/* { dg-options "-msve-vector-bits=256 -std=gnu99 -fopenmp -O2
-fdump-tree-ompexp" } */
+
+#include <arm_sve.h>
+
+#define N 256
+#define CONSTRUCT teams distribute simd
+
+svint32_t
+__attribute__ ((noinline))
+omp_target ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target CONSTRUCT
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc); /* { dg-error {'va'
referenced in target region does not have a mappable type} } */
+ }
+
+ return va;
+}
+
+svint32_t
+__attribute__ ((noinline))
+omp_target_data_map_1 ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target CONSTRUCT map(to: b, c) map(from: va) /* { dg-error {SVE
type 'svint32_t' not allowed in map clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc);
+ }
+
+ return va;
+}
+
+svint32_t
+__attribute__ ((noinline))
+omp_target_data_map_2 ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target CONSTRUCT map(to: b, c) map(from: va) /* { dg-error {SVE
type 'svint32_t' not allowed in map clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc);
+ }
+
+#pragma omp target CONSTRUCT map(to: b, c) map(tofrom: va) /* { dg-error {SVE
type 'svint32_t' not allowed in map clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, va);
+ va = svadd_s32_z (svptrue_b32 (), vc, va);
+ }
+
+ return va;
+}
+
+svint32_t
+__attribute__ ((noinline))
+omp_target_map_data_enter_exit ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target enter data map(to: b, c)
+
+#pragma omp target CONSTRUCT map(from: va) /* { dg-error {SVE type 'svint32_t'
not allowed in map clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc);
+ }
+
+#pragma omp target CONSTRUCT map(to: va) /* { dg-error {SVE type 'svint32_t'
not allowed in map clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, va);
+ va = svadd_s32_z (svptrue_b32 (), vc, va);
+ }
+
+#pragma omp target exit data map(from: va) /* { dg-error {SVE type 'svint32_t'
not allowed in map clause} } */
+
+ return va;
+}
+
+svint32_t
+__attribute__ ((noinline))
+omp_target_map_data_alloc_update ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target data map(to: b, c) map(alloc: va) /* { dg-error {SVE type
'svint32_t' not allowed in map clause} } */
+{
+#pragma omp target CONSTRUCT
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc); /* { dg-error {'va'
referenced in target region does not have a mappable type} } */
+ }
+
+/* Update va on the host from target. */
+#pragma omp target update from(va)
+
+#pragma omp target CONSTRUCT map(from: va) /* { dg-error {SVE type 'svint32_t'
not allowed in map clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, va);
+ va = svadd_s32_z (svptrue_b32 (), vc, va);
+ }
+}
+ return va;
+}
+
+int64_t __attribute__ ((noinline))
+omp_target_private ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int64_t res;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+/* Combined construct scenario: here private applies to the distribute simd
+ construct, so no error. */
+#pragma omp target CONSTRUCT private (va, vb, vc) map (to: b, c) map (from:
res)
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b);
+ vc = svld1_s32 (svptrue_b32 (), c);
+ va = svadd_s32_z (svptrue_b32 (), vb, vc);
+ res = svaddv_s32 (svptrue_b32 (), va);
+ }
+
+ return res;
+}
+
+int64_t __attribute__ ((noinline))
+omp_target_firstprivate (svbool_t vp)
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int64_t res;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target CONSTRUCT firstprivate (vp) map (to: b, c) map (from:
res)/* { dg-error {SVE type 'svbool_t' not allowed in target firstprivate
clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (vp, b); /* { dg-error {'vb' referenced in target region
does not have a mappable type} } */
+ vc = svld1_s32 (vp, c); /* { dg-error {'vc' referenced in target region
does not have a mappable type} } */
+ va = svadd_s32_z (vp, vb, vc); /* { dg-error {'va' referenced in target
region does not have a mappable type} } */
+ res = svaddv_s32 (svptrue_b32 (), va);
+ }
+
+ return res;
+}
+
diff --git
a/gcc/testsuite/gcc.target/aarch64/sve/omp/offload-teams-distribute.c
b/gcc/testsuite/gcc.target/aarch64/sve/omp/offload-teams-distribute.c
new file mode 100644
index 00000000000..6852d427866
--- /dev/null
+++ b/gcc/testsuite/gcc.target/aarch64/sve/omp/offload-teams-distribute.c
@@ -0,0 +1,236 @@
+/* { dg-do compile } */
+/* { dg-options "-msve-vector-bits=256 -std=gnu99 -fopenmp -O2
-fdump-tree-ompexp" } */
+
+#include <arm_sve.h>
+
+#define N 256
+#define CONSTRUCT teams distribute
+
+svint32_t
+__attribute__ ((noinline))
+omp_target ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target CONSTRUCT
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc); /* { dg-error {'va'
referenced in target region does not have a mappable type} } */
+ }
+
+ return va;
+}
+
+svint32_t
+__attribute__ ((noinline))
+omp_target_data_map_1 ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target CONSTRUCT map(to: b, c) map(from: va) /* { dg-error {SVE
type 'svint32_t' not allowed in map clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc);
+ }
+
+ return va;
+}
+
+svint32_t
+__attribute__ ((noinline))
+omp_target_data_map_2 ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target CONSTRUCT map(to: b, c) map(from: va) /* { dg-error {SVE
type 'svint32_t' not allowed in map clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc);
+ }
+
+#pragma omp target CONSTRUCT map(to: b, c) map(tofrom: va) /* { dg-error {SVE
type 'svint32_t' not allowed in map clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, va);
+ va = svadd_s32_z (svptrue_b32 (), vc, va);
+ }
+
+ return va;
+}
+
+svint32_t
+__attribute__ ((noinline))
+omp_target_map_data_enter_exit ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target enter data map(to: b, c)
+
+#pragma omp target CONSTRUCT map(from: va) /* { dg-error {SVE type 'svint32_t'
not allowed in map clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc);
+ }
+
+#pragma omp target CONSTRUCT map(to: va) /* { dg-error {SVE type 'svint32_t'
not allowed in map clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, va);
+ va = svadd_s32_z (svptrue_b32 (), vc, va);
+ }
+
+#pragma omp target exit data map(from: va) /* { dg-error {SVE type 'svint32_t'
not allowed in map clause} } */
+
+ return va;
+}
+
+svint32_t
+__attribute__ ((noinline))
+omp_target_map_data_alloc_update ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target data map(to: b, c) map(alloc: va) /* { dg-error {SVE type
'svint32_t' not allowed in map clause} } */
+{
+#pragma omp target CONSTRUCT
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc); /* { dg-error {'va'
referenced in target region does not have a mappable type} } */
+ }
+
+/* Update va on the host from target. */
+#pragma omp target update from(va)
+
+#pragma omp target CONSTRUCT map(from: va) /* { dg-error {SVE type 'svint32_t'
not allowed in map clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, va);
+ va = svadd_s32_z (svptrue_b32 (), vc, va);
+ }
+}
+ return va;
+}
+
+int64_t __attribute__ ((noinline))
+omp_target_private ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int64_t res;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+/* Combined construct scenario: here private applies to the teams distribute
+ construct, so no error. */
+#pragma omp target CONSTRUCT private (va, vb, vc) map (to: b, c) map (from:
res)
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b);
+ vc = svld1_s32 (svptrue_b32 (), c);
+ va = svadd_s32_z (svptrue_b32 (), vb, vc);
+ res = svaddv_s32 (svptrue_b32 (), va);
+ }
+
+ return res;
+}
+
+int64_t __attribute__ ((noinline))
+omp_target_firstprivate (svbool_t vp)
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int64_t res;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target CONSTRUCT firstprivate (vp) map (to: b, c) map (from:
res)/* { dg-error {SVE type 'svbool_t' not allowed in target firstprivate
clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (vp, b); /* { dg-error {'vb' referenced in target region
does not have a mappable type} } */
+ vc = svld1_s32 (vp, c); /* { dg-error {'vc' referenced in target region
does not have a mappable type} } */
+ va = svadd_s32_z (vp, vb, vc); /* { dg-error {'va' referenced in target
region does not have a mappable type} } */
+ res = svaddv_s32 (svptrue_b32 (), va);
+ }
+
+ return res;
+}
diff --git a/gcc/testsuite/gcc.target/aarch64/sve/omp/offload-teams-loop.c
b/gcc/testsuite/gcc.target/aarch64/sve/omp/offload-teams-loop.c
new file mode 100644
index 00000000000..aad6c47067c
--- /dev/null
+++ b/gcc/testsuite/gcc.target/aarch64/sve/omp/offload-teams-loop.c
@@ -0,0 +1,237 @@
+/* { dg-do compile } */
+/* { dg-options "-msve-vector-bits=256 -std=gnu99 -fopenmp -O2
-fdump-tree-ompexp" } */
+
+#include <arm_sve.h>
+
+#define N 256
+#define CONSTRUCT teams loop
+
+svint32_t
+__attribute__ ((noinline))
+omp_target ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target CONSTRUCT
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc); /* { dg-error {'va'
referenced in target region does not have a mappable type} } */
+ }
+
+ return va;
+}
+
+svint32_t
+__attribute__ ((noinline))
+omp_target_data_map_1 ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target CONSTRUCT map(to: b, c) map(from: va) /* { dg-error {SVE
type 'svint32_t' not allowed in map clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc);
+ }
+
+ return va;
+}
+
+svint32_t
+__attribute__ ((noinline))
+omp_target_data_map_2 ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target CONSTRUCT map(to: b, c) map(from: va) /* { dg-error {SVE
type 'svint32_t' not allowed in map clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc);
+ }
+
+#pragma omp target CONSTRUCT map(to: b, c) map(tofrom: va) /* { dg-error {SVE
type 'svint32_t' not allowed in map clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, va);
+ va = svadd_s32_z (svptrue_b32 (), vc, va);
+ }
+
+ return va;
+}
+
+svint32_t
+__attribute__ ((noinline))
+omp_target_map_data_enter_exit ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target enter data map(to: b, c)
+
+#pragma omp target CONSTRUCT map(from: va) /* { dg-error {SVE type 'svint32_t'
not allowed in map clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc);
+ }
+
+#pragma omp target CONSTRUCT map(to: va) /* { dg-error {SVE type 'svint32_t'
not allowed in map clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, va);
+ va = svadd_s32_z (svptrue_b32 (), vc, va);
+ }
+
+#pragma omp target exit data map(from: va) /* { dg-error {SVE type 'svint32_t'
not allowed in map clause} } */
+
+ return va;
+}
+
+svint32_t
+__attribute__ ((noinline))
+omp_target_map_data_alloc_update ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target data map(to: b, c) map(alloc: va) /* { dg-error {SVE type
'svint32_t' not allowed in map clause} } */
+{
+#pragma omp target CONSTRUCT
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc); /* { dg-error {'va'
referenced in target region does not have a mappable type} } */
+ }
+
+/* Update va on the host from target. */
+#pragma omp target update from(va)
+
+#pragma omp target CONSTRUCT map(from: va) /* { dg-error {SVE type 'svint32_t'
not allowed in map clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, va);
+ va = svadd_s32_z (svptrue_b32 (), vc, va);
+ }
+}
+ return va;
+}
+
+int64_t __attribute__ ((noinline))
+omp_target_private ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int64_t res;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+/* Combined construct scenario: here private applies to the teams loop
+ construct, so no error. */
+#pragma omp target CONSTRUCT private (va, vb, vc) map (to: b, c) map (from:
res)
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b);
+ vc = svld1_s32 (svptrue_b32 (), c);
+ va = svadd_s32_z (svptrue_b32 (), vb, vc);
+ res = svaddv_s32 (svptrue_b32 (), va);
+ }
+
+ return res;
+}
+
+int64_t __attribute__ ((noinline))
+omp_target_firstprivate (svbool_t vp)
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int64_t res;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target CONSTRUCT firstprivate (vp) map (to: b, c) map (from:
res)/* { dg-error {SVE type 'svbool_t' not allowed in target firstprivate
clause} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (vp, b); /* { dg-error {'vb' referenced in target region
does not have a mappable type} } */
+ vc = svld1_s32 (vp, c); /* { dg-error {'vc' referenced in target region
does not have a mappable type} } */
+ va = svadd_s32_z (vp, vb, vc); /* { dg-error {'va' referenced in target
region does not have a mappable type} } */
+ res = svaddv_s32 (svptrue_b32 (), va);
+ }
+
+ return res;
+}
+
diff --git a/gcc/testsuite/gcc.target/aarch64/sve/omp/offload-teams.c
b/gcc/testsuite/gcc.target/aarch64/sve/omp/offload-teams.c
new file mode 100644
index 00000000000..a4269108166
--- /dev/null
+++ b/gcc/testsuite/gcc.target/aarch64/sve/omp/offload-teams.c
@@ -0,0 +1,195 @@
+/* { dg-do compile } */
+/* { dg-options "-msve-vector-bits=256 -std=gnu99 -fopenmp -O2
-fdump-tree-ompexp" } */
+
+#include <arm_sve.h>
+
+#define N 256
+#define CONSTRUCT teams
+
+svint32_t
+__attribute__ ((noinline))
+omp_target ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target CONSTRUCT
+ {
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc); /* { dg-error {'va'
referenced in target region does not have a mappable type} } */
+ }
+ }
+
+ return va;
+}
+
+svint32_t
+__attribute__ ((noinline))
+omp_target_data_map_1 ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target CONSTRUCT map(to: b, c) map(from: va) /* { dg-error {SVE
type 'svint32_t' not allowed in map clause} } */
+ {
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc);
+ }
+ }
+
+ return va;
+}
+
+svint32_t
+__attribute__ ((noinline))
+omp_target_data_map_2 ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+
+#pragma omp target CONSTRUCT map(to: b, c) map(from: va) /* { dg-error {SVE
type 'svint32_t' not allowed in map clause} } */
+ {
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc);
+ }
+ }
+
+#pragma omp target CONSTRUCT map(to: b, c) map(tofrom: va) /* { dg-error {SVE
type 'svint32_t' not allowed in map clause} } */
+ {
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, va);
+ va = svadd_s32_z (svptrue_b32 (), vc, va);
+ }
+ }
+
+ return va;
+}
+
+svint32_t
+__attribute__ ((noinline))
+omp_target_map_data_enter_exit ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target enter data map(to: b, c)
+
+#pragma omp target CONSTRUCT map(from: va) /* { dg-error {SVE type 'svint32_t'
not allowed in map clause} } */
+ {
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc);
+ }
+ }
+
+#pragma omp target CONSTRUCT map(to: va) /* { dg-error {SVE type 'svint32_t'
not allowed in map clause} } */
+ {
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, va);
+ va = svadd_s32_z (svptrue_b32 (), vc, va);
+ }
+ }
+
+#pragma omp target exit data map(from: va) /* { dg-error {SVE type 'svint32_t'
not allowed in map clause} } */
+
+ return va;
+}
+
+svint32_t
+__attribute__ ((noinline))
+omp_target_map_data_alloc_update ()
+{
+
+ int a[N], b[N], c[N];
+ svint32_t va, vb, vc;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target data map(to: b, c) map(alloc: va) /* { dg-error {SVE type
'svint32_t' not allowed in map clause} } */
+{
+#pragma omp target CONSTRUCT
+ {
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, vc); /* { dg-error {'va'
referenced in target region does not have a mappable type} } */
+ }
+ }
+
+/* Update va on the host from target. */
+#pragma omp target update from(va)
+
+#pragma omp target CONSTRUCT map(from: va) /* { dg-error {SVE type 'svint32_t'
not allowed in map clause} } */
+ {
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (svptrue_b32 (), b); /* { dg-error {'vb' referenced in
target region does not have a mappable type} } */
+ vc = svld1_s32 (svptrue_b32 (), c); /* { dg-error {'vc' referenced in
target region does not have a mappable type} } */
+ va = svadd_s32_z (svptrue_b32 (), vb, va);
+ va = svadd_s32_z (svptrue_b32 (), vc, va);
+ }
+ }
+}
+ return va;
+}
diff --git a/gcc/testsuite/gcc.target/aarch64/sve/omp/target-device.c
b/gcc/testsuite/gcc.target/aarch64/sve/omp/target-device.c
new file mode 100644
index 00000000000..4c92015837f
--- /dev/null
+++ b/gcc/testsuite/gcc.target/aarch64/sve/omp/target-device.c
@@ -0,0 +1,97 @@
+/* { dg-do compile } */
+/* { dg-options "-msve-vector-bits=256 -std=gnu99 -fopenmp -O2
-fdump-tree-ompexp" } */
+
+#include <arm_sve.h>
+
+#define N 256
+
+typedef __SVInt32_t v8si __attribute__((arm_sve_vector_bits(256)));
+
+int64_t __attribute__ ((noinline))
+omp_target_device_ptr (svbool_t vp, v8si *vptr)
+{
+
+ int a[N], b[N], c[N];
+ v8si va, vb, vc;
+ int64_t res;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target data use_device_ptr (vptr) map (to: b, c) /* { dg-error
{SVE type 'v8si \*' {aka 'svint32_t
__attribute__\(\(arm_sve_vector_bits\([0-9]+\)\)\) \*'} not allowed in target
device clauses} } */
+#pragma omp target is_device_ptr (vptr) map (to: b, c) map (from: res) /* {
dg-error {SVE type 'v8si \*' {aka 'svint32_t
__attribute__\(\(arm_sve_vector_bits\(256\)\)\) \*'} not allowed in target
device clauses} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = *vptr; /* { dg-error {'vb' referenced in target region does not
have a mappable type} } */
+ vc = svld1_s32 (vp, c); /* { dg-error {'vc' referenced in target region
does not have a mappable type} } */
+ /* { dg-error {'vp' referenced in target region
does not have a mappable type} "" { target *-*-* } .-1 } */
+ va = svadd_s32_z (vp, vb, vc); /* { dg-error {'va' referenced in target
region does not have a mappable type} } */
+ res = svaddv_s32 (svptrue_b32 (), va);
+ }
+
+ return res;
+}
+
+int64_t __attribute__ ((noinline))
+omp_target_device_addr (svbool_t vp, v8si *vptr)
+{
+
+ int a[N], b[N], c[N];
+ v8si va, vb, vc;
+ int64_t res;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target data use_device_addr (vb) map (to: b, c) /* { dg-error {SVE
type 'v8si' {aka 'svint32_t __attribute__\(\(arm_sve_vector_bits\(256\)\)\)'}
not allowed in target device clauses} } */
+#pragma omp target is_device_ptr (vptr) map (to: b, c) map (from: res) /* {
dg-error {SVE type 'v8si \*' {aka 'svint32_t
__attribute__\(\(arm_sve_vector_bits\(256\)\)\) \*'} not allowed in target
device clauses} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = *vptr; /* { dg-error {'vb' referenced in target region does not
have a mappable type} } */
+ vc = svld1_s32 (vp, c); /* { dg-error {'vc' referenced in target region
does not have a mappable type} } */
+ /* { dg-error {'vp' referenced in target region
does not have a mappable type} "" { target *-*-* } .-1 } */
+ va = svadd_s32_z (vp, vb, vc); /* { dg-error {'va' referenced in target
region does not have a mappable type} } */
+ res = svaddv_s32 (svptrue_b32 (), va);
+ }
+
+ return res;
+}
+
+int64_t __attribute__ ((noinline))
+omp_target_has_device_addr (svbool_t vp, v8si *vptr)
+{
+
+ int a[N], b[N], c[N];
+ v8si va, vb, vc;
+ int64_t res;
+ int i;
+
+#pragma omp parallel for
+ for (i = 0; i < N; i++)
+ {
+ b[i] = i;
+ c[i] = i + 1;
+ }
+
+#pragma omp target data use_device_addr (vb) map (to: b, c) /* { dg-error {SVE
type 'v8si' {aka 'svint32_t __attribute__\(\(arm_sve_vector_bits\(256\)\)\)'}
not allowed in target device clauses} } */
+#pragma omp target has_device_addr (vb) map (to: b, c) map (from: res) /* {
dg-error {SVE type 'v8si' {aka 'svint32_t
__attribute__\(\(arm_sve_vector_bits\(256\)\)\)'} not allowed in target device
clauses} } */
+ for (i = 0; i < 8; i++)
+ {
+ vb = svld1_s32 (vp, b); /* { dg-error {'vp' referenced in target region
does not have a mappable type} } */
+ vc = svld1_s32 (vp, c); /* { dg-error {'vc' referenced in target region
does not have a mappable type} } */
+ va = svadd_s32_z (vp, vb, vc); /* { dg-error {'va' referenced in target
region does not have a mappable type} } */
+ res = svaddv_s32 (svptrue_b32 (), va);
+ }
+
+ return res;
+}
diff --git a/gcc/testsuite/gcc.target/aarch64/sve/omp/target-link.c
b/gcc/testsuite/gcc.target/aarch64/sve/omp/target-link.c
new file mode 100644
index 00000000000..a6e80cfd559
--- /dev/null
+++ b/gcc/testsuite/gcc.target/aarch64/sve/omp/target-link.c
@@ -0,0 +1,48 @@
+/* { dg-do compile } */
+/* { dg-options "-msve-vector-bits=256 -std=gnu99 -fopenmp -O2
-fdump-tree-ompexp" } */
+
+#include <arm_sve.h>
+
+typedef __SVInt32_t v8si __attribute__((arm_sve_vector_bits(256)));
+
+static v8si local_vec;
+#pragma omp declare target link(local_vec)
+
+v8si global_vec;
+#pragma omp declare target link(global_vec)
+
+void
+one_get_inc2_local_vec ()
+{
+ v8si res, res2, tmp;
+
+#pragma omp target map(from: res, res2) /* { dg-error {SVE type 'v8si' {aka
'svint32_t __attribute__\(\(arm_sve_vector_bits\(256\)\)\)'} not allowed in map
clause} } */
+ {
+ res = local_vec; /* { dg-error {'local_vec' referenced in target region
does not have a mappable type} } */
+ local_vec = svadd_s32_z (svptrue_b32 (), local_vec, local_vec);
+ res2 = local_vec;
+ }
+
+ tmp = svadd_s32_z (svptrue_b32 (), res, res);
+ svbool_t p = svcmpne_s32 (svptrue_b32 (), tmp, res2);
+ if (svptest_any (svptrue_b32 (), p))
+ __builtin_abort ();
+}
+
+void
+one_get_inc3_global_vec ()
+{
+ v8si res, res2, tmp;
+
+#pragma omp target map(from: res, res2) /* { dg-error {SVE type 'v8si' {aka
'svint32_t __attribute__\(\(arm_sve_vector_bits\(256\)\)\)'} not allowed in map
clause} } */
+ {
+ res = global_vec; /* { dg-error {'global_vec' referenced in target region
does not have a mappable type} } */
+ global_vec = svadd_s32_z (svptrue_b32 (), global_vec, global_vec);
+ res2 = global_vec;
+ }
+
+ tmp = svadd_s32_z (svptrue_b32 (), res, res);
+ svbool_t p = svcmpne_s32 (svptrue_b32 (), tmp, res2);
+ if (svptest_any (svptrue_b32 (), p))
+ __builtin_abort ();
+}
--
2.25.1
