Fix PR80029

In addition to resolving the memory leak involving OpenACC delare
clauses, this patch also corrects an ICE involving VLA variables as data
clause arguments used in acc declare constructs. More details can be
found here <https://gcc.gnu.org/ml/gcc-patches/2016-09/msg01630.html>.

Is this OK for trunk?

Cesar

On Wed, Mar 22, 2017 at 06:40:03AM -0700, Cesar Philippidis wrote:
> In addition to resolving the memory leak involving OpenACC delare
> clauses, this patch also corrects an ICE involving VLA variables as data
> clause arguments used in acc declare constructs. More details can be
> found here <https://gcc.gnu.org/ml/gcc-patches/2016-09/msg01630.html>.
> 
> Is this OK for trunk?
> 
> Cesar

> 2017-03-22  Cesar Philippidis  <cesar@codesourcery.com>
> 
> 	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.

Ok if testing passed.

	Jakub

On 03/22/2017 06:42 AM, Jakub Jelinek wrote:

>> 2017-03-22  Cesar Philippidis  <cesar@codesourcery.com>
>>
>> 	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.
> 
> Ok if testing passed.

Yes it did. By the way, is there a way to automate valgrind testing, or
is that something that require manual inspection?

Cesar

On Wed, Mar 22, 2017 at 06:48:09AM -0700, Cesar Philippidis wrote:
> On 03/22/2017 06:42 AM, Jakub Jelinek wrote:
> 
> >> 2017-03-22  Cesar Philippidis  <cesar@codesourcery.com>
> >>
> >> 	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.
> > 
> > Ok if testing passed.
> 
> Yes it did. By the way, is there a way to automate valgrind testing, or
> is that something that require manual inspection?

--enable-checking=valgrind if you have enough CPU time to burn (like a
weekend for a single bootstrap+regtest).

	Jakub

Hi Cesar!

On Wed, 22 Mar 2017 06:40:03 -0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
> In addition to resolving the memory leak involving OpenACC delare
> clauses, this patch also corrects an ICE involving VLA variables as data
> clause arguments used in acc declare constructs. More details can be
> found here <https://gcc.gnu.org/ml/gcc-patches/2016-09/msg01630.html>.
> 
> Is this OK for trunk?

(Got committed in r246381.)

> 	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.

> --- a/gcc/gimplify.c
> +++ b/gcc/gimplify.c
> @@ -6787,6 +6787,16 @@ device_resident_p (tree decl)
>    return false;
>  }
>  
> +/* Return true if DECL has an ACC DECLARE attribute.  */
> +
> +static bool
> +is_oacc_declared (tree decl)
> +{

Adding here a "gcc_assert (TREE_CODE (decl) != MEM_REF);", that one never
triggers (at least given the current test coverage), so...

> +  tree t = TREE_CODE (decl) == MEM_REF ? TREE_OPERAND (decl, 0) : decl;

... it would seem that this code can be simplified -- or under which
circumstances should we get a MEM_REF here?  The call from
gimplify_oacc_declare itself does the tree operand 0 indirection, and the
call from oacc_default_clause shouldn't be a MEM_REF, I would think.

> +  tree declared = lookup_attribute ("oacc declare target", DECL_ATTRIBUTES (t));
> +  return declared != NULL_TREE;
> +}

> @@ -6887,6 +6897,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))
> @@ -6917,7 +6928,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

Isn't this new "declared" logic doing a thing very similar to the
existing "on_device" logic?  Should that be combined/cleaned up?

Why doesn't the same "declared"/"on_device" logic also apply to
ORT_ACC_KERNELS?


Grüße
 Thomas


> @@ -7346,6 +7357,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:
> @@ -9231,18 +9243,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)
> @@ -9258,7 +9278,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/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 <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;
> +}