Hi Jakub,

On 2021/9/18 12:11 AM, Jakub Jelinek wrote:
>> @@ -3496,7 +3509,8 @@ static match
>>   match_omp (gfc_exec_op op, const omp_mask mask)
>>   {
>>     gfc_omp_clauses *c;
>> -  if (gfc_match_omp_clauses (&c, mask) != MATCH_YES)
>> +  if (gfc_match_omp_clauses (&c, mask, true, true, false,
>> +			     (op == EXEC_OMP_TARGET)) != MATCH_YES)
> 
> The ()s around op == EXEC_OMP_TARGET are unnecessary.

Fixed.

>> --- a/gcc/fortran/trans-openmp.c
>> +++ b/gcc/fortran/trans-openmp.c
>> @@ -6391,12 +6391,17 @@ gfc_trans_omp_task (gfc_code *code)
>>   static tree
>>   gfc_trans_omp_taskgroup (gfc_code *code)
>>   {
>> +  stmtblock_t block;
>> +  gfc_start_block (&block);
>>     tree body = gfc_trans_code (code->block->next);
>>     tree stmt = make_node (OMP_TASKGROUP);
>>     TREE_TYPE (stmt) = void_type_node;
>>     OMP_TASKGROUP_BODY (stmt) = body;
>> -  OMP_TASKGROUP_CLAUSES (stmt) = NULL_TREE;
>> -  return stmt;
>> +  OMP_TASKGROUP_CLAUSES (stmt) = gfc_trans_omp_clauses (&block,
>> +							code->ext.omp_clauses,
>> +							code->loc);
>> +  gfc_add_expr_to_block (&block, stmt);
> 
> If this was missing, then I'm afraid we lack a lot of testsuite coverage for
> Fortran task reductions.  It doesn't need to be covered in this patch, but would be
> good to cover it incrementally.  Because the above means nothing with
> taskgroup with task_reduction clause(s) could work properly at runtime.

Actually, the testcases do somewhat exercise taskgroup task_reductions, but like you
said, only lightly.

>> --- a/gcc/omp-low.c
>> +++ b/gcc/omp-low.c
>> @@ -1317,9 +1317,13 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
>>   	  if (is_omp_target (ctx->stmt))
>>   	    {
>>   	      tree at = decl;
>> +	      omp_context *scan_ctx = ctx;
>>   	      if (ctx->outer)
>> -		scan_omp_op (&at, ctx->outer);
>> -	      tree nt = omp_copy_decl_1 (at, ctx);
>> +		{
>> +		  scan_omp_op (&at, ctx->outer);
>> +		  scan_ctx = ctx->outer;
>> +		}
>> +	      tree nt = omp_copy_decl_1 (at, scan_ctx);
>>   	      splay_tree_insert (ctx->field_map,
>>   				 (splay_tree_key) &DECL_CONTEXT (decl),
>>   				 (splay_tree_value) nt);
> 
> You're right that the var remembered with &DECL_CONTEXT (whatever) key is
> used outside of the target construct rather than inside of it.
> So, if ctx->outer is non-NULL, it seems right to create the var in that
> outer context.  But, if ctx->outer is NULL, which can happen if the
> target construct is orphaned, consider e.g.
> extern int &x;
> extern int &y;
> 
> void
> foo ()
> {
>    #pragma omp target in_reduction (+: x, y)
>    {
>      x = x + 8;
>      y = y + 16;
>    }
> }
> 
> void
> bar ()
> {
>    #pragma omp taskgroup task_reduction (+: x, y)
>    foo ();
> }
> then those artificial decls (copies of x and y) should appear
> to be at the function scope and not inside of the target region.
> 
> Therefore, I wonder if omp_copy_decl_2 shouldn't do the
>    DECL_CONTEXT (copy) = current_function_decl;
>    DECL_CHAIN (copy) = ctx->block_vars;
>    ctx->block_vars = copy;
> (the last one can be moved next to the others) only if ctx != NULL
> and otherwise call gimple_add_tmp_var (copy); instead
> and then just call omp_copy_decl_1 at that spot with unconditional
> ctx->outer.

I see what you mean. I tried gimple_add_tmp_var but didn't work due to a
!DECL_SEEN_IN_BIND_EXPR_P() assert fail, but record_vars() appears to work.

> Also, this isn't the only place that should have such a change,
> there is also
>                    if (ctx->outer)
>                      scan_omp_op (&at, ctx->outer);
>                    tree nt = omp_copy_decl_1 (at, ctx);
>                    splay_tree_insert (ctx->field_map,
>                                       (splay_tree_key) &DECL_CONTEXT (t),
>                                       (splay_tree_value) nt);
> a few lines above this and I'd expect that it should be (at, ctx->outer)
> as well.

Fixed.

>> @@ -1339,7 +1343,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
>>   	      if (!is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx)))
>>   		{
>>   		  by_ref = use_pointer_for_field (decl, ctx);
>> -		  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IN_REDUCTION)
>> +		  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IN_REDUCTION
>> +		      && !splay_tree_lookup (ctx->field_map,
>> +					     (splay_tree_key) decl))
>>   		    install_var_field (decl, by_ref, 3, ctx);
>>   		}
>>   	      install_var_local (decl, ctx);
> 
> When exactly do you need this?  It doesn't trigger on the new libgomp
> testcase...

I remember there was a testcase with triggered an ICE without this, but for some
reason can't find it anymore. I don't have any more evidence this is needed, so
removed now.

>> --- /dev/null
>> +++ b/libgomp/testsuite/libgomp.fortran/target-in-reduction-1.f90
>> @@ -0,0 +1,33 @@
>> +! { dg-do run }
>> +
>> +subroutine foo (x, y)
...
>> +  if (x .ne. 11) stop 1
>> +  if (y .ne. 21) stop 2
>> +
>> +end program main
> 
> Again, something that can be dealt incrementally, but the
> testsuite coverage of
> https://gcc.gnu.org/pipermail/gcc-patches/2021-June/573600.html
> was larger than this.  Would be nice e.g. to cover both scalar vars
> and array sections/arrays, parameters passed by reference as in the
> above testcase, but also something that isn't a reference (either a local
> variable or dummy parameter with VALUE, etc.
> 
> 	Jakub

I have expanded target-in-reduction-1.f90 to cover local variables and
VALUE passed parameters. Array sections in reductions appear to be still
not supported by the Fortran FE in general (Tobias plans to work on that later).

I also added another target-in-reduction-2.f90 testcase that tests the "orphaned"
case in Fortran, where the task/target-in_reduction is in another separate subroutine.

Tested without regressions on trunk, is this okay to commit?

Thanks,
Chung-Lin

2021-10-19  Chung-Lin Tang  <cltang@codesourcery.com>

gcc/fortran/ChangeLog:

	* openmp.c (gfc_match_omp_clause_reduction): Add 'openmp_target' default
	false parameter. Add 'always,tofrom' map for OMP_LIST_IN_REDUCTION case.
	(gfc_match_omp_clauses): Add 'openmp_target' default false parameter,
	adjust call to gfc_match_omp_clause_reduction.
	(match_omp): Adjust call to gfc_match_omp_clauses
	* trans-openmp.c (gfc_trans_omp_taskgroup): Add call to
	gfc_match_omp_clause, create and return block.

gcc/ChangeLog:

	* omp-low.c (omp_copy_decl_2): For !ctx, use record_vars to add new copy
	as local variable.
	(scan_sharing_clauses): Place copy of OMP_CLAUSE_IN_REDUCTION decl in
	ctx->outer instead of ctx.

gcc/testsuite/ChangeLog:

	* gfortran.dg/gomp/reduction4.f90: Adjust omp target in_reduction' scan
	pattern.

libgomp/ChangeLog:

	* testsuite/libgomp.fortran/target-in-reduction-1.f90: New test.
	* testsuite/libgomp.fortran/target-in-reduction-2.f90: New test.