Commit 7ba8651e by Cesar Philippidis

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
parent bf634d1c
t2017-03-22 Thomas Preud'homme <thomas.preudhomme@arm.com> 2017-03-22 Cesar Philippidis <cesar@codesourcery.com>
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 <thomas.preudhomme@arm.com>
PR target/80082 PR target/80082
* config/arm/arm-isa.h (isa_bit_lpae): New feature bit. * config/arm/arm-isa.h (isa_bit_lpae): New feature bit.
......
...@@ -6786,6 +6786,16 @@ device_resident_p (tree decl) ...@@ -6786,6 +6786,16 @@ device_resident_p (tree decl)
return false; 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 /* Determine outer default flags for DECL mentioned in an OMP region
but not declared in an enclosing clause. but not declared in an enclosing clause.
...@@ -6886,6 +6896,7 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags) ...@@ -6886,6 +6896,7 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
{ {
const char *rkind; const char *rkind;
bool on_device = false; bool on_device = false;
bool declared = is_oacc_declared (decl);
tree type = TREE_TYPE (decl); tree type = TREE_TYPE (decl);
if (lang_hooks.decls.omp_privatize_by_reference (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) ...@@ -6916,7 +6927,7 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
case ORT_ACC_PARALLEL: 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'. */ /* Aggregates default to 'present_or_copy'. */
flags |= GOVD_MAP; flags |= GOVD_MAP;
else else
...@@ -7345,6 +7356,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, ...@@ -7345,6 +7356,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
case OMP_TARGET_DATA: case OMP_TARGET_DATA:
case OMP_TARGET_ENTER_DATA: case OMP_TARGET_ENTER_DATA:
case OMP_TARGET_EXIT_DATA: case OMP_TARGET_EXIT_DATA:
case OACC_DECLARE:
case OACC_HOST_DATA: case OACC_HOST_DATA:
ctx->target_firstprivatize_array_bases = true; ctx->target_firstprivatize_array_bases = true;
default: default:
...@@ -9230,18 +9242,26 @@ gimplify_oacc_declare (tree *expr_p, gimple_seq *pre_p) ...@@ -9230,18 +9242,26 @@ gimplify_oacc_declare (tree *expr_p, gimple_seq *pre_p)
{ {
tree expr = *expr_p; tree expr = *expr_p;
gomp_target *stmt; gomp_target *stmt;
tree clauses, t; tree clauses, t, decl;
clauses = OACC_DECLARE_CLAUSES (expr); clauses = OACC_DECLARE_CLAUSES (expr);
gimplify_scan_omp_clauses (&clauses, pre_p, ORT_TARGET_DATA, OACC_DECLARE); 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)) 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) 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) if (VAR_P (decl)
&& !is_global_var (decl) && !is_global_var (decl)
...@@ -9257,6 +9277,7 @@ gimplify_oacc_declare (tree *expr_p, gimple_seq *pre_p) ...@@ -9257,6 +9277,7 @@ gimplify_oacc_declare (tree *expr_p, gimple_seq *pre_p)
} }
} }
if (gimplify_omp_ctxp)
omp_add_variable (gimplify_omp_ctxp, decl, GOVD_SEEN); omp_add_variable (gimplify_omp_ctxp, decl, GOVD_SEEN);
} }
......
2017-03-22 Cesar Philippidis <cesar@codesourcery.com>
PR c++/80029
* testsuite/libgomp.oacc-c-c++-common/declare-vla.c: New test.
2017-03-08 Jakub Jelinek <jakub@redhat.com> 2017-03-08 Jakub Jelinek <jakub@redhat.com>
PR c/79940 PR c/79940
......
/* Verify that acc declare accept VLA variables. */
#include <assert.h>
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;
}
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment