This is the mail archive of the gcc-patches@gcc.gnu.org mailing list for the GCC project.


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]
Other format: [Raw text]

Re: [OpenACC] Enable firstprivate OpenACC reductions


On Wed, Sep 05, 2018 at 01:08:21PM -0700, Cesar Philippidis wrote:
> 2018-XX-YY  Cesar Philippidis  <cesar@codesourcery.com>
> 	    Chung-Lin Tang  <cltang@codesourcery.com>
> 
> 	gcc/
> 	* gimplify.c (omp_add_variable): Enable firstprivate reduction
> 	variables.
> 
> 	gcc/testsuite/
> 	* c-c++-common/goacc/reduction-8.c: New test.
> 
> 	libgomp/
> 	* testsuite/libgomp.oacc-c-c++-common/privatize-reduction-1.c: New
> 	test.
> 	* testsuite/libgomp.oacc-c-c++-common/privatize-reduction-2.c: New
> 	test.

This is all in ORT_ACC guarded code and I don't see anything wrong with it,
so if Thomas is ok with it, it is ok for trunk.

> --- a/gcc/gimplify.c
> +++ b/gcc/gimplify.c
> @@ -6823,20 +6823,27 @@ omp_add_variable (struct gimplify_omp_ctx *ctx, tree decl, unsigned int flags)
>    else
>      splay_tree_insert (ctx->variables, (splay_tree_key)decl, flags);
>  
> -  /* For reductions clauses in OpenACC loop directives, by default create a
> -     copy clause on the enclosing parallel construct for carrying back the
> -     results.  */
> +  /* For OpenACC loop directives, when a reduction clause is placed on
> +     the outermost acc loop within an acc parallel or kernels
> +     construct, it must have an implied copy data mapping. E.g.
> +
> +       #pragma acc parallel
> +	 {
> +	   #pragma acc loop reduction (+:sum)
> +
> +     a copy clause for sum should be added on the enclosing parallel
> +     construct for carrying back the results.  */
>    if (ctx->region_type == ORT_ACC && (flags & GOVD_REDUCTION))
>      {
>        struct gimplify_omp_ctx *outer_ctx = ctx->outer_context;
> -      while (outer_ctx)
> +      if (outer_ctx)
>  	{
>  	  n = splay_tree_lookup (outer_ctx->variables, (splay_tree_key)decl);
>  	  if (n != NULL)
>  	    {
>  	      /* Ignore local variables and explicitly declared clauses.  */
>  	      if (n->value & (GOVD_LOCAL | GOVD_EXPLICIT))
> -		break;
> +		;
>  	      else if (outer_ctx->region_type == ORT_ACC_KERNELS)
>  		{
>  		  /* According to the OpenACC spec, such a reduction variable
> @@ -6856,9 +6863,7 @@ omp_add_variable (struct gimplify_omp_ctx *ctx, tree decl, unsigned int flags)
>  	    {
>  	      splay_tree_insert (outer_ctx->variables, (splay_tree_key)decl,
>  				 GOVD_MAP | GOVD_SEEN);
> -	      break;
>  	    }
> -	  outer_ctx = outer_ctx->outer_context;
>  	}
>      }
>  }

	Jakub


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]