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