Re: [gomp4] Handle deviceptr from an outer directive

[Thomas Schwinge <thomas at codesourcery dot com>]

Hi Jim!

On Tue, 7 Jul 2015 10:19:39 -0500, James Norris <jnorris@codesourcery.com> wrote:
> This patch fixes an issue where the deviceptr clause in an outer
> directive was being ignored during implicit variable definition
> on a nested directive.

> Committed to gomp-4_0-branch.

> --- a/gcc/gimplify.c
> +++ b/gcc/gimplify.c

I'm sorry, have not yet tried very hard; but I can't claim to understand
the logic here -- why is the OpenACC deviceptr clause so special?  :-|

> @@ -116,6 +116,9 @@ enum gimplify_omp_var_data
>    /* Gang-local OpenACC variable.  */
>    GOVD_GANGLOCAL = (1 << 16),
>  
> +  /* OpenACC deviceptr clause.  */
> +  GOVD_USE_DEVPTR = (1 << 17),
> +
>    GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
>  			   | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
>  			   | GOVD_LOCAL)
> @@ -6274,7 +6277,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
>  		}
>  	      break;
>  	    }
> +
>  	  flags = GOVD_MAP | GOVD_EXPLICIT;
> +	  if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_DEVICEPTR)
> +	    flags |= GOVD_USE_DEVPTR;
>  	  goto do_add;
>  
>  	case OMP_CLAUSE_DEPEND:
> @@ -6662,6 +6668,7 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
>  			       : (flags & GOVD_FORCE_MAP
>  				  ? GOMP_MAP_FORCE_TOFROM
>  				  : GOMP_MAP_TOFROM));
> +
>        if (DECL_SIZE (decl)
>  	  && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
>  	{
> @@ -6687,7 +6694,17 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
>  	  OMP_CLAUSE_CHAIN (clause) = nc;
>  	}
>        else
> -	OMP_CLAUSE_SIZE (clause) = DECL_SIZE_UNIT (decl);
> +	{
> +	  if (gimplify_omp_ctxp->outer_context)
> +	    {
> +	      struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp->outer_context;
> +	      splay_tree_node on
> +		    = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
> +	      if (on && (on->value & GOVD_USE_DEVPTR))
> +	        OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_FORCE_PRESENT);
> +	    }
> +	  OMP_CLAUSE_SIZE (clause) = DECL_SIZE_UNIT (decl);
> +	}
>      }
>    if (code == OMP_CLAUSE_FIRSTPRIVATE && (flags & GOVD_LASTPRIVATE) != 0)
>      {

The patch that you committed (r225518) also includes a test case
(thanks!) as follows:

    --- /dev/null
    +++ gcc/testsuite/c-c++-common/goacc/deviceptr-4.c
    @@ -0,0 +1,12 @@
    +/* { dg-additional-options "-fdump-tree-gimple" } */
    +
    +void
    +subr (int *a)
    +{
    +#pragma acc data deviceptr (a)
    +#pragma acc parallel
    +  a[0] += 1.0;
    +}
    +
    +/* { dg-final { scan-tree-dump-times "#pragma omp target oacc_parallel.*map\\(force_present:a \\\[len: 8\\\]\\)" 1 "gimple" } } */
    +/* { dg-final { cleanup-tree-dump "gimple" } } */

That len: 8 is obviously valid only for 64-bit configurations, so will
cause a FAIL on anything else.


GrÃÃe,
 Thomas

Attachment: signature.asc
Description: PGP signature