Hi Thomas!

On 07/09/2015 03:51 AM, Thomas Schwinge wrote:
> 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?  :-|

Because no data movement should be initiated on behalf of the user.

Prior to the fix, a map (tofrom) was being inserted when the variable
in question was implied within a nested directive. This is wrong.
No movement should occur as the memory for the variable is already
present on the target. Therefore, the use of 'present' is the correct
specification.


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

Fixed and committed to gomp-4_0-branch.

Thank you, thank you,
Jim