Warn about inconsistent OpenACC nested reduction clauses

	OpenACC (cf. OpenACC 2.7, section 2.9.11. "reduction clause";
	this was first clarified by OpenACC 2.6) requires that, if a
	variable is used in reduction clauses on two nested loops, then
	there must be reduction clauses for that variable on all loops
	that are nested in between the two loops and all these reduction
	clauses must use the same operator.
	This commit introduces a check for that property which reports
	warnings if it is violated.

	2019-11-06  Gergö Barany  <gergo@codesourcery.com>
	            Frederik Harwath  <frederik@codesourcery.com>
	            Thomas Schwinge  <thomas@codesourcery.com>

	gcc/
	* omp-low.c (struct omp_context): New fields
	local_reduction_clauses, outer_reduction_clauses.
	(new_omp_context): Initialize these.
	(scan_sharing_clauses): Record reduction clauses on OpenACC constructs.
	(scan_omp_for): Check reduction clauses for incorrect nesting.
	gcc/testsuite/
	* c-c++-common/goacc/nested-reductions-warn.c: New test.
	* c-c++-common/goacc/nested-reductions.c: New test.
	* gfortran.dg/goacc/nested-reductions-warn.f90: New test.
	* gfortran.dg/goacc/nested-reductions.f90: New test.
	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-1.c:
	Add expected warnings about missing reduction clauses.
	* testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-2.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-3.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-4.c:
	Likewise.

	Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>

From-SVN: r277875

gcc/ChangeLog

+2019-11-06  Gergö Barany  <gergo@codesourcery.com>
+	    Frederik Harwath  <frederik@codesourcery.com>
+	    Thomas Schwinge  <thomas@codesourcery.com>
+
+	* omp-low.c (struct omp_context): New fields
+	local_reduction_clauses, outer_reduction_clauses.
+	(new_omp_context): Initialize these.
+	(scan_sharing_clauses): Record reduction clauses on OpenACC constructs.
+	(scan_omp_for): Check reduction clauses for incorrect nesting.
+	
 2019-11-06  Jakub Jelinek  <jakub@redhat.com>
 
 	PR inline-asm/92352

gcc/omp-low.c

@@ -128,6 +128,12 @@ struct omp_context
      corresponding tracking loop iteration variables.  */
   hash_map<tree, tree> *lastprivate_conditional_map;
 
+  /* A tree_list of the reduction clauses in this context.  */
+  tree local_reduction_clauses;
+
+  /* A tree_list of the reduction clauses in outer contexts.  */
+  tree outer_reduction_clauses;
+
   /* Nesting depth of this context.  Used to beautify error messages re
      invalid gotos.  The outermost ctx is depth 1, with depth 0 being
      reserved for the main body of the function.  */
@@ -910,6 +916,8 @@ new_omp_context (gimple *stmt, omp_context *outer_ctx)
       ctx->outer = outer_ctx;
       ctx->cb = outer_ctx->cb;
       ctx->cb.block = NULL;
+      ctx->local_reduction_clauses = NULL;
+      ctx->outer_reduction_clauses = ctx->outer_reduction_clauses;
       ctx->depth = outer_ctx->depth + 1;
     }
   else
@@ -925,6 +933,8 @@ new_omp_context (gimple *stmt, omp_context *outer_ctx)
       ctx->cb.transform_call_graph_edges = CB_CGE_MOVE;
       ctx->cb.adjust_array_error_bounds = true;
       ctx->cb.dont_remap_vla_if_no_change = true;
+      ctx->local_reduction_clauses = NULL;
+      ctx->outer_reduction_clauses = NULL;
       ctx->depth = 1;
     }
 
@@ -1139,6 +1149,11 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	  goto do_private;
 
 	case OMP_CLAUSE_REDUCTION:
+	  if (is_oacc_parallel (ctx) || is_oacc_kernels (ctx))
+	    ctx->local_reduction_clauses
+	      = tree_cons (NULL, c, ctx->local_reduction_clauses);
+	  /* FALLTHRU */
+
 	case OMP_CLAUSE_IN_REDUCTION:
 	  decl = OMP_CLAUSE_DECL (c);
 	  if (TREE_CODE (decl) == MEM_REF)
@@ -2423,6 +2438,88 @@ scan_omp_for (gomp_for *stmt, omp_context *outer_ctx)
 	  gimple_omp_for_set_clauses (stmt, clauses);
 	  check_oacc_kernel_gwv (stmt, ctx);
 	}
+
+      /* Collect all variables named in reductions on this loop.  Ensure
+	 that, if this loop has a reduction on some variable v, and there is
+	 a reduction on v somewhere in an outer context, then there is a
+	 reduction on v on all intervening loops as well.  */
+      tree local_reduction_clauses = NULL;
+      for (tree c = gimple_omp_for_clauses (stmt); c; c = OMP_CLAUSE_CHAIN (c))
+	{
+	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
+	    local_reduction_clauses
+	      = tree_cons (NULL, c, local_reduction_clauses);
+	}
+      if (ctx->outer_reduction_clauses == NULL && ctx->outer != NULL)
+	ctx->outer_reduction_clauses
+	  = chainon (unshare_expr (ctx->outer->local_reduction_clauses),
+		     ctx->outer->outer_reduction_clauses);
+      tree outer_reduction_clauses = ctx->outer_reduction_clauses;
+      tree local_iter = local_reduction_clauses;
+      for (; local_iter; local_iter = TREE_CHAIN (local_iter))
+	{
+	  tree local_clause = TREE_VALUE (local_iter);
+	  tree local_var = OMP_CLAUSE_DECL (local_clause);
+	  tree_code local_op = OMP_CLAUSE_REDUCTION_CODE (local_clause);
+	  bool have_outer_reduction = false;
+	  tree ctx_iter = outer_reduction_clauses;
+	  for (; ctx_iter; ctx_iter = TREE_CHAIN (ctx_iter))
+	    {
+	      tree outer_clause = TREE_VALUE (ctx_iter);
+	      tree outer_var = OMP_CLAUSE_DECL (outer_clause);
+	      tree_code outer_op = OMP_CLAUSE_REDUCTION_CODE (outer_clause);
+	      if (outer_var == local_var && outer_op != local_op)
+		{
+		  warning_at (gimple_location (stmt), 0,
+			      "conflicting reduction operations for %qE",
+			      local_var);
+		  inform (OMP_CLAUSE_LOCATION (outer_clause),
+			  "location of the previous reduction for %qE",
+			  outer_var);
+		}
+	      if (outer_var == local_var)
+		{
+		  have_outer_reduction = true;
+		  break;
+		}
+	    }
+	  if (have_outer_reduction)
+	    {
+	      /* There is a reduction on outer_var both on this loop and on
+		 some enclosing loop.  Walk up the context tree until such a
+		 loop with a reduction on outer_var is found, and complain
+		 about all intervening loops that do not have such a
+		 reduction.  */
+	      struct omp_context *curr_loop = ctx->outer;
+	      bool found = false;
+	      while (curr_loop != NULL)
+		{
+		  tree curr_iter = curr_loop->local_reduction_clauses;
+		  for (; curr_iter; curr_iter = TREE_CHAIN (curr_iter))
+		    {
+		      tree curr_clause = TREE_VALUE (curr_iter);
+		      tree curr_var = OMP_CLAUSE_DECL (curr_clause);
+		      if (curr_var == local_var)
+			{
+			  found = true;
+			  break;
+			}
+		    }
+		  if (!found)
+		    warning_at (gimple_location (curr_loop->stmt), 0,
+				"nested loop in reduction needs "
+				"reduction clause for %qE",
+				local_var);
+		  else
+		    break;
+		  curr_loop = curr_loop->outer;
+		}
+	    }
+	}
+      ctx->local_reduction_clauses = local_reduction_clauses;
+      ctx->outer_reduction_clauses
+	= chainon (unshare_expr (ctx->local_reduction_clauses),
+		   ctx->outer_reduction_clauses);
     }
 
   scan_sharing_clauses (clauses, ctx);

gcc/testsuite/ChangeLog

+2019-11-06  Gergö Barany  <gergo@codesourcery.com>
+	    Frederik Harwath  <frederik@codesourcery.com>
+	    Thomas Schwinge  <thomas@codesourcery.com>
+
+	* c-c++-common/goacc/nested-reductions-warn.c: New test.
+	* c-c++-common/goacc/nested-reductions.c: New test.
+	* gfortran.dg/goacc/nested-reductions-warn.f90: New test.
+	* gfortran.dg/goacc/nested-reductions.f90: New test.
+
 2019-11-06  Jakub Jelinek  <jakub@redhat.com>
 
 	PR inline-asm/92352

libgomp/ChangeLog

+2019-11-06  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-1.c:
+	Add expected warnings about missing reduction clauses.
+	* testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-2.c:
+	Likewise.
+	* testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-3.c:
+	Likewise.
+	* testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-4.c:
+	Likewise.
+
 2019-11-04  Tobias Burnus  <tobias@codesourcery.com>
 
 	* testsuite/libgomp.fortran/pr66199-1.f90: Remove

libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-1.c

@@ -15,7 +15,7 @@ main (int argc, char *argv[])
   #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
     reduction(+:res) copy(res)
   {
-    #pragma acc loop gang
+    #pragma acc loop gang /* { dg-warning "nested loop in reduction needs reduction clause for 'res'" "TODO" } */
     for (j = 0; j < 32; j++)
       {
 	#pragma acc loop worker reduction(+:res)

libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-2.c

@@ -14,7 +14,7 @@ main (int argc, char *argv[])
   #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
 		       reduction(^:res)
   {
-    #pragma acc loop gang
+    #pragma acc loop gang /* { dg-warning "nested loop in reduction needs reduction clause for 'res'" "TODO" } */
     for (j = 0; j < 32; j++)
       {
 	#pragma acc loop worker vector reduction(^:res)

libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-3.c

@@ -16,7 +16,7 @@ main (int argc, char *argv[])
   #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
     reduction(+:res) copy(res)
   {
-    #pragma acc loop gang
+    #pragma acc loop gang /* { dg-warning "nested loop in reduction needs reduction clause for 'res'" "TODO" } */
     for (j = 0; j < 32; j++)
       {
 	#pragma acc loop worker vector reduction(+:res)

libgomp/testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-4.c

@@ -16,7 +16,7 @@ main (int argc, char *argv[])
   #pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
     reduction(+:res) reduction(max:mres) copy(res, mres)
   {
-    #pragma acc loop gang
+    #pragma acc loop gang /* { dg-warning "nested loop in reduction needs reduction clause for 'm\?res'" "TODO" } */
     for (j = 0; j < 32; j++)
       {
 	#pragma acc loop worker vector reduction(+:res) reduction(max:mres)