From 7ba8651ed253b29ee17a0d735cc74eac1b826492 Mon Sep 17 00:00:00 2001 From: Cesar Philippidis Date: Wed, 22 Mar 2017 06:52:10 -0700 Subject: [PATCH] re PR c++/80029 (valgrind error in new_omp_context(omp_region_type) (gimplify.c:400)) PR c++/80029 gcc/ * gimplify.c (is_oacc_declared): New function. (oacc_default_clause): Use it to set default flags for acc declared variables inside parallel regions. (gimplify_scan_omp_clauses): Strip firstprivate pointers for acc declared variables. (gimplify_oacc_declare): Gimplify the declare clauses. Add the declare attribute to any decl as necessary. libgomp/ * testsuite/libgomp.oacc-c-c++-common/declare-vla.c: New test. From-SVN: r246381 --- gcc/ChangeLog | 13 ++++++++- gcc/gimplify.c | 31 ++++++++++++++++++---- libgomp/ChangeLog | 5 ++++ .../libgomp.oacc-c-c++-common/declare-vla.c | 25 +++++++++++++++++ 4 files changed, 68 insertions(+), 6 deletions(-) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 18d00c2..f7202cf 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,4 +1,15 @@ -t2017-03-22 Thomas Preud'homme +2017-03-22 Cesar Philippidis + + PR c++/80029 + * gimplify.c (is_oacc_declared): New function. + (oacc_default_clause): Use it to set default flags for acc declared + variables inside parallel regions. + (gimplify_scan_omp_clauses): Strip firstprivate pointers for acc + declared variables. + (gimplify_oacc_declare): Gimplify the declare clauses. Add the + declare attribute to any decl as necessary. + +2017-03-22 Thomas Preud'homme PR target/80082 * config/arm/arm-isa.h (isa_bit_lpae): New feature bit. diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 06f984c..5658d0a 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -6786,6 +6786,16 @@ device_resident_p (tree decl) return false; } +/* Return true if DECL has an ACC DECLARE attribute. */ + +static bool +is_oacc_declared (tree decl) +{ + tree t = TREE_CODE (decl) == MEM_REF ? TREE_OPERAND (decl, 0) : decl; + tree declared = lookup_attribute ("oacc declare target", DECL_ATTRIBUTES (t)); + return declared != NULL_TREE; +} + /* Determine outer default flags for DECL mentioned in an OMP region but not declared in an enclosing clause. @@ -6886,6 +6896,7 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags) { const char *rkind; bool on_device = false; + bool declared = is_oacc_declared (decl); tree type = TREE_TYPE (decl); if (lang_hooks.decls.omp_privatize_by_reference (decl)) @@ -6916,7 +6927,7 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags) case ORT_ACC_PARALLEL: { - if (on_device || AGGREGATE_TYPE_P (type)) + if (on_device || AGGREGATE_TYPE_P (type) || declared) /* Aggregates default to 'present_or_copy'. */ flags |= GOVD_MAP; else @@ -7345,6 +7356,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, case OMP_TARGET_DATA: case OMP_TARGET_ENTER_DATA: case OMP_TARGET_EXIT_DATA: + case OACC_DECLARE: case OACC_HOST_DATA: ctx->target_firstprivatize_array_bases = true; default: @@ -9230,18 +9242,26 @@ gimplify_oacc_declare (tree *expr_p, gimple_seq *pre_p) { tree expr = *expr_p; gomp_target *stmt; - tree clauses, t; + tree clauses, t, decl; clauses = OACC_DECLARE_CLAUSES (expr); gimplify_scan_omp_clauses (&clauses, pre_p, ORT_TARGET_DATA, OACC_DECLARE); + gimplify_adjust_omp_clauses (pre_p, NULL, &clauses, OACC_DECLARE); for (t = clauses; t; t = OMP_CLAUSE_CHAIN (t)) { - tree decl = OMP_CLAUSE_DECL (t); + decl = OMP_CLAUSE_DECL (t); if (TREE_CODE (decl) == MEM_REF) - continue; + decl = TREE_OPERAND (decl, 0); + + if (VAR_P (decl) && !is_oacc_declared (decl)) + { + tree attr = get_identifier ("oacc declare target"); + DECL_ATTRIBUTES (decl) = tree_cons (attr, NULL_TREE, + DECL_ATTRIBUTES (decl)); + } if (VAR_P (decl) && !is_global_var (decl) @@ -9257,7 +9277,8 @@ gimplify_oacc_declare (tree *expr_p, gimple_seq *pre_p) } } - omp_add_variable (gimplify_omp_ctxp, decl, GOVD_SEEN); + if (gimplify_omp_ctxp) + omp_add_variable (gimplify_omp_ctxp, decl, GOVD_SEEN); } stmt = gimple_build_omp_target (NULL, GF_OMP_TARGET_KIND_OACC_DECLARE, diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index f63f028..74f50e0 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,8 @@ +2017-03-22 Cesar Philippidis + + PR c++/80029 + * testsuite/libgomp.oacc-c-c++-common/declare-vla.c: New test. + 2017-03-08 Jakub Jelinek PR c/79940 diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c new file mode 100644 index 0000000..3ea148e --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c @@ -0,0 +1,25 @@ +/* Verify that acc declare accept VLA variables. */ + +#include + +int +main () +{ + int N = 1000; + int i, A[N]; +#pragma acc declare copy(A) + + for (i = 0; i < N; i++) + A[i] = -i; + +#pragma acc kernels + for (i = 0; i < N; i++) + A[i] = i; + +#pragma acc update host(A) + + for (i = 0; i < N; i++) + assert (A[i] == i); + + return 0; +} -- 2.7.4