c-typeck.c (c_finish_omp_clauses): Mark OpenACC reduction arguments as addressable...

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

diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog
index d2a8947d1b39700801521792eb0d37b260659d77..1262c82cedb2ff97e887cc6f2b4bab8e97aab4ed 100644 (file)
--- a/gcc/c/ChangeLog
+++ b/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

diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index 1520c20c653f42fdc6221ebda436d1d2c61f8ff8..0ff28f2ec2810c50f51adce76e4c43536277f7ff 100644 (file)
--- a/gcc/c/c-typeck.c
+++ b/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);

diff --git a/gcc/cp/ChangeLog b/gcc/cp/ChangeLog
index 69d3da17c9ece2ca26d3d10c0408d1ae29675ffe..bec708c9f817efa3b7e44199629e2c24cc606e80 100644 (file)
--- a/gcc/cp/ChangeLog
+++ b/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

diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 8e682c537955456dd69a54f71c3ee2754f795fce..536509127e4fbd5d4689772872ba00c31b47cfdc 100644 (file)
--- a/gcc/cp/semantics.c
+++ b/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;

diff --git a/gcc/fortran/ChangeLog b/gcc/fortran/ChangeLog
index 2a9c9ceab934ab845cd0d6781c8895c23530b699..8c04089f8d02d6ff637da458a99750d2b1c4b35d 100644 (file)
--- a/gcc/fortran/ChangeLog
+++ b/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

diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index d3276f9331f7dfb58fef6b2610cc36e2252691f1..ab07fe45be9e1951a1631a1eb018328608e341bb 100644 (file)
--- a/gcc/fortran/trans-openmp.c
+++ b/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;

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-8.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-8.c
new file mode 100644 (file)
index 0000000..d250605
--- /dev/null
+++ b/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;
+}

diff --git a/libgomp/testsuite/libgomp.oacc-fortran/reduction-8.f90 b/libgomp/testsuite/libgomp.oacc-fortran/reduction-8.f90
new file mode 100644 (file)
index 0000000..5cf4681
--- /dev/null
+++ b/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