c-typeck.c (c_finish_omp_clauses): Mark OpenACC reduction arguments as addressable when async clause exists.

2016-06-03  Chung-Lin Tang  <cltang@codesourcery.com>

	c/
	* c-typeck.c (c_finish_omp_clauses): Mark OpenACC reduction
	arguments as addressable when async clause exists.

	cp/
	* semantics.c (finish_omp_clauses): Mark OpenACC reduction
	arguments as addressable when async clause exists.

	fortran/
	* trans-openmp.c (gfc_trans_omp_reduction_list): Add mark_addressable
	bool parameter, set reduction clause DECLs as addressable when true.
	(gfc_trans_omp_clauses): Pass clauses->async to
	gfc_trans_omp_reduction_list, add comment describing OpenACC situation.

	libgomp/
	* testsuite/libgomp.oacc-fortran/reduction-8.f90: New testcase.
	* testsuite/libgomp.oacc-c-c++-common/reduction-8.c: New testcase.

From-SVN: r237070

gcc/c/ChangeLog

@@ -1,3 +1,8 @@
+2016-06-03  Chung-Lin Tang  <cltang@codesourcery.com>
+
+	* c-typeck.c (c_finish_omp_clauses): Mark OpenACC reduction
+	arguments as addressable when async clause exists.
+
 2016-05-30  Jakub Jelinek  <jakub@redhat.com>
 
 	PR c++/71349

gcc/c/c-typeck.c

@@ -12529,6 +12529,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
   tree *nowait_clause = NULL;
   bool ordered_seen = false;
   tree schedule_clause = NULL_TREE;
+  bool oacc_async = false;
 
   bitmap_obstack_initialize (NULL);
   bitmap_initialize (&generic_head, &bitmap_default_obstack);
@@ -12539,6 +12540,14 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
   bitmap_initialize (&map_field_head, &bitmap_default_obstack);
   bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack);
 
+  if (ort & C_ORT_ACC)
+    for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_ASYNC)
+	{
+	  oacc_async = true;
+	  break;
+	}
+	  
   for (pc = &clauses, c = clauses; c ; c = *pc)
     {
       bool remove = false;
@@ -12575,6 +12584,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	      remove = true;
 	      break;
 	    }
+	  if (oacc_async)
+	    c_mark_addressable (t);
 	  type = TREE_TYPE (t);
 	  if (TREE_CODE (t) == MEM_REF)
 	    type = TREE_TYPE (type);

gcc/cp/ChangeLog

@@ -1,3 +1,8 @@
+2016-06-03  Chung-Lin Tang  <cltang@codesourcery.com>
+
+	* semantics.c (finish_omp_clauses): Mark OpenACC reduction
+	arguments as addressable when async clause exists.
+
 2016-06-02  Jan Hubicka  <jh@suse.cz>
 
 	* cp-gimplify.c (genericize_continue_stmt): Force addition of

gcc/cp/semantics.c

@@ -5775,6 +5775,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
   bool branch_seen = false;
   bool copyprivate_seen = false;
   bool ordered_seen = false;
+  bool oacc_async = false;
 
   bitmap_obstack_initialize (NULL);
   bitmap_initialize (&generic_head, &bitmap_default_obstack);
@@ -5785,6 +5786,14 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
   bitmap_initialize (&map_field_head, &bitmap_default_obstack);
   bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack);
 
+  if (ort & C_ORT_ACC)
+    for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_ASYNC)
+	{
+	  oacc_async = true;
+	  break;
+	}
+
   for (pc = &clauses, c = clauses; c ; c = *pc)
     {
       bool remove = false;
@@ -5828,6 +5837,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		t = n;
 	      goto check_dup_generic_t;
 	    }
+	  if (oacc_async)
+	    cxx_mark_addressable (t);
 	  goto check_dup_generic;
 	case OMP_CLAUSE_COPYPRIVATE:
 	  copyprivate_seen = true;

gcc/fortran/ChangeLog

@@ -1,3 +1,10 @@
+2016-06-03  Chung-Lin Tang  <cltang@codesourcery.com>
+
+	* trans-openmp.c (gfc_trans_omp_reduction_list): Add mark_addressable
+	bool parameter, set reduction clause DECLs as addressable when true.
+	(gfc_trans_omp_clauses): Pass clauses->async to
+	gfc_trans_omp_reduction_list, add comment describing OpenACC situation.
+
 2016-06-01  Jerry DeLisle  <jvdelisle@gcc.gnu.org>
 
 	PR fortran/52393

gcc/fortran/trans-openmp.c

@@ -1646,7 +1646,7 @@ gfc_trans_omp_array_reduction_or_udr (tree c, gfc_omp_namelist *n, locus where)
 
 static tree
 gfc_trans_omp_reduction_list (gfc_omp_namelist *namelist, tree list,
-			      locus where)
+			      locus where, bool mark_addressable)
 {
   for (; namelist != NULL; namelist = namelist->next)
     if (namelist->sym->attr.referenced)
@@ -1657,6 +1657,8 @@ gfc_trans_omp_reduction_list (gfc_omp_namelist *namelist, tree list,
 	    tree node = build_omp_clause (where.lb->location,
 					  OMP_CLAUSE_REDUCTION);
 	    OMP_CLAUSE_DECL (node) = t;
+	    if (mark_addressable)
+	      TREE_ADDRESSABLE (t) = 1;
 	    switch (namelist->u.reduction_op)
 	      {
 	      case OMP_REDUCTION_PLUS:
@@ -1747,7 +1749,10 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
       switch (list)
 	{
 	case OMP_LIST_REDUCTION:
-	  omp_clauses = gfc_trans_omp_reduction_list (n, omp_clauses, where);
+	  /* An OpenACC async clause indicates the need to set reduction
+	     arguments addressable, to allow asynchronous copy-out.  */
+	  omp_clauses = gfc_trans_omp_reduction_list (n, omp_clauses, where,
+						      clauses->async);
 	  break;
 	case OMP_LIST_PRIVATE:
 	  clause_code = OMP_CLAUSE_PRIVATE;

libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-8.c

@@ -0,0 +1,30 @@
+const int n = 100;
+    
+// Check async over parallel construct with reduction
+    
+int
+async_sum (int c)
+{
+  int s = 0;
+    
+#pragma acc parallel loop num_gangs (10) gang reduction (+:s) async
+  for (int i = 0; i < n; i++)
+    s += i+c;
+    
+#pragma acc wait
+  return s;
+}
+    
+int
+main()
+{
+  int result = 0;
+    
+  for (int i = 0; i < n; i++)
+    result += i+1;
+    
+  if (async_sum (1) != result)
+    __builtin_abort ();
+    
+  return 0;
+}

libgomp/testsuite/libgomp.oacc-fortran/reduction-8.f90

@@ -0,0 +1,41 @@
+! { dg-do run }
+
+program reduction
+  implicit none
+  integer, parameter    :: n = 100
+  integer               :: i, h1, h2, s1, s2, a1, a2
+
+  h1 = 0
+  h2 = 0
+  do i = 1, n
+     h1 = h1 + 1
+     h2 = h2 + 2
+  end do
+
+  s1 = 0
+  s2 = 0
+  !$acc parallel loop reduction(+:s1, s2)
+  do i = 1, n
+     s1 = s1 + 1
+     s2 = s2 + 2
+  end do
+  !$acc end parallel loop
+
+  a1 = 0
+  a2 = 0
+  !$acc parallel loop reduction(+:a1, a2) async(1)
+  do i = 1, n
+     a1 = a1 + 1
+     a2 = a2 + 2
+  end do
+  !$acc end parallel loop
+
+  if (h1 .ne. s1) call abort ()
+  if (h2 .ne. s2) call abort ()
+
+  !$acc wait(1)
+
+  if (h1 .ne. a1) call abort ()
+  if (h2 .ne. a2) call abort ()
+
+end program reduction