Re: [PATCH] [gomp4] Initial support of OpenACC loop directive in C front-end.

[Thomas Schwinge <thomas at codesourcery dot com>]

Hi!

On Tue, 5 Aug 2014 08:22:50 -0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
> On 07/29/2014 02:07 AM, Thomas Schwinge wrote:
> 
> > On Thu, 20 Mar 2014 15:42:48 +0100, I wrote:
> >> On Tue, 18 Mar 2014 14:50:44 +0100, I wrote:
> >>> On Tue, 18 Mar 2014 16:37:24 +0400, Ilmir Usmanov <i.usmanov@samsung.com> wrote:
> >>>> This patch introduces support of OpenACC loop directive (and combined 
> >>>> directives) in C front-end up to GENERIC. Currently no clause is allowed.
> >>>
> >>> Thanks!  I had worked on a simpler patch, not yet dealing with combined
> >>> clauses.  Also, I have some work for the GIMPLE level, namely building on
> >>> GIMPLE_OMP_FOR, adding a new GF_OMP_FOR_KIND_OACC_LOOP.  I'll post this
> >>> soon.
> >>
> >> Here are the patches, committed in r208702..4 to gomp-4_0-branch.
> > 
> > Cesar, I hope I'm not confusing things here, but I remember that you once
> > pointed out that in Fortran OpenACC, we have to explicitly specify the
> > loop iteration variable in a private clause, whereas the OpenACC
> > specification says this needs to happen automatically (predetermined data
> > attribute).

> > In light of this, please also review whether the following
> > gimplify_omp_for changes of mine (when I originally added the OACC_LOOP
> > support) are correct.  Evidently, this code still has TODO markers in it;
> > this was well before you added preliminary support for the private
> > clause.  That is, should we use GOVD_PRIVATE also for OACC_LOOP?
> 
> OMP has both a private and a shared clause. And those neither of those
> clauses can share variables, i.e. an index variable cannot be both
> private and shared. By default, index variables are private in OMP, so
> the gimplification pass makes those variables explicitly private so that
> later passes can check for errors.

> > Why does this nevertheless currently work for C for loops without any
> > private clause for the loop iteration variable?  Maybe because in C, the
> > loop iteration variable ends up in a register and that is not shared
> > between threads?  I do see a private clause being added during
> > gimplification for the OpenMP C test case, but not for OpenACC -- which I
> > assume is precisely due to the code I'm quoting below.
> 
> OMP lowering pass makes the index variables local to the parallel
> function/kernel. So each thread effectively has its private copy of the
> index variable.

I have now committed the following to gomp-4_0-branch in r217433:

commit 8562c72ad30e73dbd77b9d4170f71a395357aa33
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date:   Wed Nov 12 17:49:59 2014 +0000

    OpenACC loop construct: predetermined data attributes for loop variables.
    
    	gcc/
    	* gimplify.c (gimplify_omp_for): Eliminate govd_private; always
    	use GOVD_PRIVATE.
    	gcc/testsuite/
    	* c-c++-common/goacc/loop-private-1.c: New file.
    	libgomp/testsuite/
    	* testsuite/libgomp.oacc-c-c++-common/collapse-4.c: New file.
    	* testsuite/libgomp.oacc-c/collapse-4.c: Turn into an execution
    	test, check result.
    	* testsuite/libgomp.oacc-c-c++-common/data-2.c: Update.
    	* testsuite/libgomp.oacc-c-c++-common/data-3.c: Likewise.
    
    git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@217433 138bc75d-0d04-0410-961f-82ee72b054a4
---
 gcc/ChangeLog.gomp                                   |  3 +++
 gcc/gimplify.c                                       | 13 ++++---------
 gcc/testsuite/ChangeLog.gomp                         |  4 ++++
 gcc/testsuite/c-c++-common/goacc/loop-private-1.c    | 14 ++++++++++++++
 libgomp/ChangeLog.gomp                               |  8 ++++++++
 .../collapse-4.c                                     | 12 +++++++-----
 libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c | 20 ++++++++++----------
 libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c | 20 ++++++++++----------
 libgomp/testsuite/libgomp.oacc-c/collapse-4.c        |  6 ++++--
 9 files changed, 64 insertions(+), 36 deletions(-)

diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index 2008542..0269746 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,5 +1,8 @@
 2014-11-12  Thomas Schwinge  <thomas@codesourcery.com>
 
+	* gimplify.c (gimplify_omp_for): Eliminate govd_private; always
+	use GOVD_PRIVATE.
+
 	* omp-low.c (scan_oacc_offload, expand_oacc_offload)
 	(lower_oacc_offload): Merge into scan_omp_target,
 	expand_omp_target, lower_omp_target, respectively.  Update all
diff --git gcc/gimplify.c gcc/gimplify.c
index 5aeb726..233ac56 100644
--- gcc/gimplify.c
+++ gcc/gimplify.c
@@ -6843,7 +6843,6 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
   gimple_seq for_body, for_pre_body;
   int i;
   bool simd;
-  enum gimplify_omp_var_data govd_private;
   enum omp_region_type ort;
   bitmap has_decl_expr = NULL;
 
@@ -6855,18 +6854,15 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
     case CILK_FOR:
     case OMP_DISTRIBUTE:
       simd = false;
-      govd_private = GOVD_PRIVATE;
       ort = ORT_WORKSHARE;
       break;
     case OMP_SIMD:
     case CILK_SIMD:
       simd = true;
-      govd_private = GOVD_PRIVATE;
       ort = ORT_SIMD;
       break;
     case OACC_LOOP:
       simd = false;
-      govd_private = /* TODO */ GOVD_LOCAL;
       ort = /* TODO */ ORT_WORKSHARE;
       break;
     default:
@@ -6928,7 +6924,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
       gcc_assert (INTEGRAL_TYPE_P (TREE_TYPE (decl))
 		  || POINTER_TYPE_P (TREE_TYPE (decl)));
 
-      /* Make sure the iteration variable is some kind of private.  */
+      /* Make sure the iteration variable is private.  */
       tree c = NULL_TREE;
       tree c2 = NULL_TREE;
       if (orig_for_stmt != for_stmt)
@@ -6957,7 +6953,6 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
 	    }
 	  else
 	    {
-	      gcc_assert (govd_private == GOVD_PRIVATE);
 	      bool lastprivate
 		= (!has_decl_expr
 		   || !bitmap_bit_p (has_decl_expr, DECL_UID (decl)));
@@ -7001,7 +6996,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
       else if (omp_is_private (gimplify_omp_ctxp, decl, 0))
 	omp_notice_variable (gimplify_omp_ctxp, decl, true);
       else
-	omp_add_variable (gimplify_omp_ctxp, decl, govd_private | GOVD_SEEN);
+	omp_add_variable (gimplify_omp_ctxp, decl, GOVD_PRIVATE | GOVD_SEEN);
 
       /* If DECL is not a gimple register, create a temporary variable to act
 	 as an iteration counter.  This is valid, since DECL cannot be
@@ -7036,7 +7031,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
 	    }
 	  else
 	    omp_add_variable (gimplify_omp_ctxp, var,
-			      govd_private | GOVD_SEEN);
+			      GOVD_PRIVATE | GOVD_SEEN);
 	}
       else
 	var = decl;
@@ -7193,7 +7188,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
 	t = TREE_VEC_ELT (OMP_FOR_INIT (for_stmt), i);
 	decl = TREE_OPERAND (t, 0);
 	var = create_tmp_var (TREE_TYPE (decl), get_name (decl));
-	omp_add_variable (gimplify_omp_ctxp, var, govd_private | GOVD_SEEN);
+	omp_add_variable (gimplify_omp_ctxp, var, GOVD_PRIVATE | GOVD_SEEN);
 	TREE_OPERAND (t, 0) = var;
 	t = TREE_VEC_ELT (OMP_FOR_INCR (for_stmt), i);
 	TREE_OPERAND (t, 1) = copy_node (TREE_OPERAND (t, 1));
diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp
index f8bacc3..a979d31 100644
--- gcc/testsuite/ChangeLog.gomp
+++ gcc/testsuite/ChangeLog.gomp
@@ -1,3 +1,7 @@
+2014-11-12  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* c-c++-common/goacc/loop-private-1.c: New file.
+
 2014-11-07  Thomas Schwinge  <thomas@codesourcery.com>
 	    James Norris  <jnorris@codesourcery.com>
 
diff --git gcc/testsuite/c-c++-common/goacc/loop-private-1.c gcc/testsuite/c-c++-common/goacc/loop-private-1.c
new file mode 100644
index 0000000..a54edb2
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/loop-private-1.c
@@ -0,0 +1,14 @@
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+void
+f (int i, int j)
+{
+#pragma acc kernels
+#pragma acc loop collapse(2)
+  for (i = 0; i < 20; ++i)
+    for (j = 0; j < 25; ++j)
+      ;
+}
+
+/* { dg-final { scan-tree-dump-times "#pragma acc loop collapse\\(2\\) private\\(j\\) private\\(i\\)" 1 "gimple" } } */
+/* { dg-final { cleanup-tree-dump "gimple" } } */
diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp
index 9f562d5..36beb59 100644
--- libgomp/ChangeLog.gomp
+++ libgomp/ChangeLog.gomp
@@ -1,3 +1,11 @@
+2014-11-12  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* testsuite/libgomp.oacc-c-c++-common/collapse-4.c: New file.
+	* testsuite/libgomp.oacc-c/collapse-4.c: Turn into an execution
+	test, check result.
+	* testsuite/libgomp.oacc-c-c++-common/data-2.c: Update.
+	* testsuite/libgomp.oacc-c-c++-common/data-3.c: Likewise.
+
 2014-11-07  Thomas Schwinge  <thomas@codesourcery.com>
 
 	* plugin-nvptx.c: Remove.
diff --git libgomp/testsuite/libgomp.oacc-c/collapse-4.c libgomp/testsuite/libgomp.oacc-c-c++-common/collapse-4.c
similarity index 56%
copy from libgomp/testsuite/libgomp.oacc-c/collapse-4.c
copy to libgomp/testsuite/libgomp.oacc-c-c++-common/collapse-4.c
index 7df0de2..f3538a6 100644
--- libgomp/testsuite/libgomp.oacc-c/collapse-4.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/collapse-4.c
@@ -1,6 +1,4 @@
-/* TODO: change to a run test once libgomp supports all data clauses.  */
-/* { dg-do compile } */
-/* { dg-options "-O2 -std=c99" } */
+/* See also ../libgomp.oacc-c/collapse-4.c.  */
 
 #include <string.h>
 
@@ -9,17 +7,21 @@ main (void)
 {
   int l = 0;
   int b[3][3];
+  int i, j;
 
   memset (b, '\0', sizeof (b));
 
 #pragma acc parallel copy(b[0:3][0:3]) copy(l)
     {
 #pragma acc loop collapse(2) reduction(+:l)
-	for (int i = 0; i < 2; i++)
-	  for (int j = 0; j < 2; j++)
+	for (i = 0; i < 2; i++)
+	  for (j = 0; j < 2; j++)
 	    if (b[i][j] != 16)
 		  l += 1;
     }
 
+  if (l != 2 * 2)
+    __builtin_abort();
+
   return 0;
 }
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c
index b990ade..3e02ffa 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c
@@ -28,8 +28,8 @@ main (int argc, char **argv)
 #pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async
 #pragma acc parallel async wait
 #pragma acc loop
-  for (int ii = 0; ii < N; ii++)
-    b[ii] = a[ii];
+  for (i = 0; i < N; i++)
+    b[i] = a[i];
 
 #pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) wait async
 #pragma acc wait
@@ -52,8 +52,8 @@ main (int argc, char **argv)
 #pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async (1)
 #pragma acc parallel async (1)
 #pragma acc loop
-  for (int ii = 0; ii < N; ii++)
-    b[ii] = a[ii];
+  for (i = 0; i < N; i++)
+    b[i] = a[i];
 
 #pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) wait (1) async (1)
 #pragma acc wait (1)
@@ -79,18 +79,18 @@ main (int argc, char **argv)
 
 #pragma acc parallel async (1) wait (1)
 #pragma acc loop
-  for (int ii = 0; ii < N; ii++)
-    b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
+  for (i = 0; i < N; i++)
+    b[i] = (a[i] * a[i] * a[i]) / a[i];
 
 #pragma acc parallel async (2) wait (1)
 #pragma acc loop
-  for (int ii = 0; ii < N; ii++)
-    c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
+  for (i = 0; i < N; i++)
+    c[i] = (a[i] + a[i] + a[i] + a[i]) / a[i];
 
 #pragma acc parallel async (3) wait (1)
 #pragma acc loop
-  for (int ii = 0; ii < N; ii++)
-    d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
+  for (i = 0; i < N; i++)
+    d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i];
 
 #pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) copyout (d[0:N]) wait (1, 2, 3) async (1)
 #pragma acc wait (1)
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c
index f8f1b3b..44787dd 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c
@@ -28,8 +28,8 @@ main (int argc, char **argv)
 #pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async
 #pragma acc parallel async wait
 #pragma acc loop
-  for (int ii = 0; ii < N; ii++)
-    b[ii] = a[ii];
+  for (i = 0; i < N; i++)
+    b[i] = a[i];
 
 #pragma acc update host (a[0:N], b[0:N]) async wait
 #pragma acc wait
@@ -52,8 +52,8 @@ main (int argc, char **argv)
 #pragma acc update device (a[0:N], b[0:N]) async (1)
 #pragma acc parallel async (1)
 #pragma acc loop
-  for (int ii = 0; ii < N; ii++)
-    b[ii] = a[ii];
+  for (i = 0; i < N; i++)
+    b[i] = a[i];
 
 #pragma acc update host (a[0:N], b[0:N]) async (1) wait (1)
 #pragma acc wait (1)
@@ -81,18 +81,18 @@ main (int argc, char **argv)
 
 #pragma acc parallel async (1) wait (1,2)
 #pragma acc loop
-  for (int ii = 0; ii < N; ii++)
-    b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
+  for (i = 0; i < N; i++)
+    b[i] = (a[i] * a[i] * a[i]) / a[i];
 
 #pragma acc parallel async (2) wait (1,3)
 #pragma acc loop
-  for (int ii = 0; ii < N; ii++)
-    c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
+  for (i = 0; i < N; i++)
+    c[i] = (a[i] + a[i] + a[i] + a[i]) / a[i];
 
 #pragma acc parallel async (3) wait (1,3)
 #pragma acc loop
-  for (int ii = 0; ii < N; ii++)
-    d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
+  for (i = 0; i < N; i++)
+    d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i];
 
 #pragma acc update host (a[0:N], b[0:N], c[0:N], d[0:N]) async (1) wait (1,2,3)
 #pragma acc wait (1)
diff --git libgomp/testsuite/libgomp.oacc-c/collapse-4.c libgomp/testsuite/libgomp.oacc-c/collapse-4.c
index 7df0de2..17663ca 100644
--- libgomp/testsuite/libgomp.oacc-c/collapse-4.c
+++ libgomp/testsuite/libgomp.oacc-c/collapse-4.c
@@ -1,5 +1,4 @@
-/* TODO: change to a run test once libgomp supports all data clauses.  */
-/* { dg-do compile } */
+/* See also ../libgomp.oacc-c-c++-common/collapse-4.c.  */
 /* { dg-options "-O2 -std=c99" } */
 
 #include <string.h>
@@ -21,5 +20,8 @@ main (void)
 		  l += 1;
     }
 
+  if (l != 2 * 2)
+    __builtin_abort();
+
   return 0;
 }


GrÃÃe,
 Thomas

Attachment: signature.asc
Description: PGP signature