| Message ID | 20240527050626.3769230-4-tejas.belagod@arm.com |
|---|---|
| State | New |
| Headers | show |
| Series | AArch64/OpenMP: Test SVE ACLE types with various OpenMP constructs. | expand |
Tejas Belagod <tejas.belagod@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.corallo@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 (); > +}
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 (); +}