[OpenMP,Fortran] Support in_reduction for Fortran

  Hi Jakub, and Fortran folks,
this patch does the required adjustments to let 'in_reduction' work for Fortran.
Not just for the target directive actually, task directive is also working after
this patch.

There is a little bit of adjustment in omp-low.c:scan_sharing_clauses:
RTL expand of the copy of the OMP_CLAUSE_IN_REDUCTION decl was failing
for Fortran by-reference arguments, which seems to work after placing them
under the outer ctx (when it exists). This also now needs checking the field_map
for existence of the field before inserting.

Tested without regressions on mainline trunk, is this okay?

(testing for devel/omp/gcc-11 is in progress)

Thanks,
Chung-Lin

2021-09-17  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 (scan_sharing_clauses): Place in_reduction copy of variable
	in outer ctx if if exists. Check if non-existent in field_map before
	installing OMP_CLAUSE_IN_REDUCTION decl.

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.
  

On Fri, Sep 17, 2021 at 07:57:38PM +0800, Chung-Lin Tang wrote:
> 2021-09-17  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 (scan_sharing_clauses): Place in_reduction copy of variable
> 	in outer ctx if if exists. Check if non-existent in field_map before
> 	installing OMP_CLAUSE_IN_REDUCTION decl.
> 
> 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.

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

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

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

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.

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

> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.fortran/target-in-reduction-1.f90
> @@ -0,0 +1,33 @@
> +! { dg-do run }
> +
> +subroutine foo (x, y)
> +  integer :: x, y
> +
> +  !$omp taskgroup task_reduction (+: x, y)
> +
> +  !$omp target in_reduction (+: x, y)
> +  x = x + 8
> +  y = y + 16
> +  !$omp end target
> +
> +  !$omp task in_reduction (+: x, y)
> +  x = x + 2
> +  y = y + 4
> +  !$omp end task
> +
> +  !$omp end taskgroup
> +
> +end subroutine foo
> +
> +program main
> +  integer :: x, y
> +
> +  x = 1
> +  y = 1
> +
> +  call 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