Tejas Belagod <tejas.bela...@arm.com> writes: > 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))))
Could you say in more detail why we check for zero or one levels of pointer indirection but not for more? Also, was there a reason for checking builtin_type_p rather than sizeless_type_p? Things like svbool_t remain sizeless even for -msve-vector-bits=128 etc., so sizeless_type_p would still cover that case. But arm_sve_vector_bits makes it possible to define fixed-length vector types that are treated for ABI & ACLE purposes like SVE types. I don't think those should be treated differently from normal vectors by omp, since the size is fixed by the attribute (and types with different attributes are distinct). Thanks, Richard > + 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 (); > +}