Re: [Patch][Fortran] OpenACC – permit common blocks in some clauses

[Tobias Burnus]

Hi Thomas,

On 10/25/19 10:43 AM, Thomas Schwinge wrote:
> OK for trunk, with the following few small items considered.

Committed as Rev. 277451 – after a fresh bootstrap and regtesting.

Changes:
* I have now a new test case 
libgomp/testsuite/libgomp.oacc-fortran/common-block-3.f90 which looks at 
omplower.
* In the compile-time *{2,3} test case, there is now also a 'enter data' 
and 'update host/self/device' test.
* the libgomp tests have a 'dg-do run'.
* I modified the code in gimplify.c as proposed.


Regarding the new test case: Without the gcc/gimplify.c changes, one has 
(see last item before child fn):

 Â Â Â  #pragma omp target oacc_parallel map(tofrom:a [len: 400]) 
map(tofrom:b [len: 400]) map(tofrom:c [len: 4]) map(tofrom:block [len: 
812]) [child fn …
 Â Â Â  #pragma omp target oacc_kernels map(force_tofrom:i [len: 4]) 
map(tofrom:y [len: 400]) map(tofrom:x [len: 400]) 
map(tofrom:kernel_block [len: 804]) map(force_tofrom:c [len: 4]) 
map(tofrom:block [len: 812])  [child fn …

With the changes of gcc/gimplify.c, one has:

 Â Â Â  #pragma omp target oacc_parallel map(tofrom:a [len: 400]) 
map(tofrom:b [len: 400]) map(tofrom:c [len: 4]) [child fn …
 Â Â Â  #pragma omp target oacc_kernels map(force_tofrom:i [len: 4]) 
map(tofrom:y [len: 400]) map(tofrom:x [len: 400]) map(force_tofrom:c 
[len: 4])  [child fn …


And without gimplify.c, the added run-tests indeed fail with:
libgomp: Trying to map into device [0x407100..0x407294) object when 
[0x407100..0x407290) is already mapped


Tobias

PS:
> Or, would it be easy to add an OpenACC 'kernels' test case that otherwise
> faild (at run time, say, with aforementioned duplicate mapping errors, or
> would contain "strange"/duplicate/conflicting mapping items in the
> '-fdump-tree-gimple' dump)?

See new test case and result for the current tests.

Additionally, I have applied:

> Wouldn't it be clearer if that latter one were written as follows:
>      if (DECL_HAS_VALUE_EXPR_P (decl))
>        {
>          if (ctx->region_type & ORT_ACC)
>            /* For OpenACC, defer expansion of value to avoid transfering
>               privatized common block data instead of im-/explicitly transfered
>               variables which are in common blocks.  */
>            ;
>          else
>            {
>              tree value = get_base_address (DECL_VALUE_EXPR (decl));
>      
>              if (value && DECL_P (value) && DECL_THREAD_LOCAL_P (value))
>                return omp_notice_threadprivate_variable (ctx, decl, value);
>            }
>        }
>
>> @@ -7353,7 +7374,9 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
>>     n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
>>     if ((ctx->region_type & ORT_TARGET) != 0)
>>       {
>> -      ret = lang_hooks.decls.omp_disregard_value_expr (decl, true);
>> +      /* For OpenACC, as remarked above, defer expansion.  */
>> +      shared = !(ctx->region_type & ORT_ACC);
>> +      ret = lang_hooks.decls.omp_disregard_value_expr (decl, shared);
> Also more explicit, easier to read:
>
>      if (ctx->region_type & ORT_ACC)
>        /* For OpenACC, as remarked above, defer expansion.  */
>        shared = false;
>      else
>        shared = true;
>
>> @@ -7521,6 +7544,9 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
>>       }
>>   
>>     shared = ((flags | n->value) & GOVD_SHARED) != 0;
>> +  /* For OpenACC, cf. remark above regaring common blocks.  */
>> +  if (ctx->region_type & ORT_ACC)
>> +    shared = false;
>>     ret = lang_hooks.decls.omp_disregard_value_expr (decl, shared);
> And again:
>
>      if (ctx->region_type & ORT_ACC)
>        /* For OpenACC, cf. remark above regaring common blocks.  */
>        shared = false;
>      else
>        shared = ((flags | n->value) & GOVD_SHARED) != 0;
>
> (In all three cases, using an easy 'if (ctx->region_type & ORT_ACC)' to
> point out the special case.)
>
> It's still some kind of voodoo to me -- but at least, you've now also
> reviewed this, and it's now documented what's going on.


And changed the test case based on:

>> +  !$acc exit data delete(/blockA/, /blockB/, e, v)
> I note there is one single 'exit data' test, but no 'enter data'.
>
> Also, 'update' is missing, to test the 'device' and 'self'/'host' clauses.
>
>> +  !$acc data deviceptr(/blockA/, /blockB/, e, v) ! { dg-error "Syntax error in OpenMP variable list" }
>> +  !$acc end data ! { dg-error "Unexpected ..ACC END DATA statement" }
>> +
>> +  !$acc data deviceptr(/blockA/, /blockB/, e, v) ! { dg-error "Syntax error in OpenMP variable list" }
>> +  !$acc end data ! { dg-error "Unexpected ..ACC END DATA statement" }
>> +end program test
> Is there a reason for the duplicated 'deviceptr' testing?
>
> Move 'data deviceptr' up a little bit, next to the other 'data' construct
> testing?
>
>> --- /dev/null
>> +++ b/gcc/testsuite/gfortran.dg/goacc/common-block-2.f90
> Similarly.
-------------- next part --------------
A non-text attachment was scrubbed...
Name: acc-common-committed.diff
Type: text/x-patch
Size: 27768 bytes
Desc: not available
URL: <http://gcc.gnu.org/pipermail/gcc-patches/attachments/20191025/929669a3/attachment.bin>