Re: [gomp4.1] Depend clause support for offloading

[Jakub Jelinek <jakub@redhat.com>]

On Wed, Sep 02, 2015 at 05:58:54PM +0200, Jakub Jelinek wrote:
> Here is the start of the async offloading support I've talked about,
> but nowait is not supported on the library side yet, only depend clause
> (and for that I haven't added a testcase yet).

Added testcase revealed two (small) issues, here is a fix for that together
with the testcase.

BTW, unless we want to add (at least now) support for running tasks in
between sending offloading target requests for memory allocation or data
movement and the offloading target signalizing their completion (supposedly
we'd better then be able to perform something like writev, merge as many
requests as possible into one metarequest and then await the completion of
it), I think at least for now we can ignore nowait on
target {update,{enter,exit} data} if depend clause is not also present
(on the library side).

I'll try to work on target {update,{enter,exit} data} nowait depend next
(in that case we need to copy the arrays and create some gomp_task).

2015-09-03  Jakub Jelinek  <jakub@redhat.com>

	* omp-low.c (lower_depend_clauses): Set TREE_ADDRESSABLE on array.
	(lower_omp_target): Use gimple_omp_target_clauses_ptr instead of
	gimple_omp_task_clauses_ptr.

	* testsuite/libgomp.c/target-25.c: New test.

--- gcc/omp-low.c.jj	2015-09-02 15:13:13.000000000 +0200
+++ gcc/omp-low.c	2015-09-03 15:24:15.153716381 +0200
@@ -12975,6 +12975,7 @@ lower_depend_clauses (tree *pclauses, gi
 	}
   tree type = build_array_type_nelts (ptr_type_node, n_in + n_out + 2);
   tree array = create_tmp_var (type);
+  TREE_ADDRESSABLE (array) = 1;
   tree r = build4 (ARRAY_REF, ptr_type_node, array, size_int (0), NULL_TREE,
 		   NULL_TREE);
   g = gimple_build_assign (r, build_int_cst (ptr_type_node, n_in + n_out));
@@ -13182,7 +13183,7 @@ lower_omp_target (gimple_stmt_iterator *
     {
       push_gimplify_context ();
       dep_bind = gimple_build_bind (NULL, NULL, make_node (BLOCK));
-      lower_depend_clauses (gimple_omp_task_clauses_ptr (stmt),
+      lower_depend_clauses (gimple_omp_target_clauses_ptr (stmt),
 			    &dep_ilist, &dep_olist);
     }
 
--- libgomp/testsuite/libgomp.c/target-25.c.jj	2015-09-03 15:02:34.130651945 +0200
+++ libgomp/testsuite/libgomp.c/target-25.c	2015-09-03 15:49:52.077362256 +0200
@@ -0,0 +1,84 @@
+#include <stdlib.h>
+#include <unistd.h>
+
+int
+main ()
+{
+  int x = 0, y = 0, z = 0, s = 11, t = 12, u = 13, w = 7, err;
+  #pragma omp parallel
+  #pragma omp single
+  {
+    #pragma omp task depend(in: x)
+    {
+      usleep (5000);
+      x = 1;
+    }
+    #pragma omp task depend(in: x)
+    {
+      usleep (6000);
+      y = 2;
+    }
+    #pragma omp task depend(out: z)
+    {
+      usleep (7000);
+      z = 3;
+    }
+    #pragma omp target map(tofrom: x) firstprivate (y) depend(inout: x, z)
+    err = (x != 1 || y != 2 || z != 3);
+    if (err)
+      abort ();
+    #pragma omp task depend(in: x)
+    {
+      usleep (5000);
+      x = 4;
+    }
+    #pragma omp task depend(in: x)
+    {
+      usleep (4000);
+      y = 5;
+    }
+    #pragma omp task depend(in: z)
+    {
+      usleep (3000);
+      z = 6;
+    }
+    #pragma omp target enter data nowait map (to: w)
+    #pragma omp target enter data depend (inout: x, z) map (to: x, y, z)
+    #pragma omp target map (alloc: x, y, z)
+    {
+      err = (x != 4 || y != 5 || z != 6);
+      x = 7;
+      y = 8;
+      z = 9;
+    }
+    if (err)
+      abort ();
+    #pragma omp taskwait
+    #pragma omp target map (alloc: w)
+    {
+      err = w != 7;
+      w = 17;
+    }
+    if (err)
+      abort (); 
+    #pragma omp task depend(in: x)
+    {
+      usleep (2000);
+      s = 14;
+    }
+    #pragma omp task depend(in: x)
+    {
+      usleep (3000);
+      t = 15;
+    }
+    #pragma omp task depend(in: z)
+    {
+      usleep (4000);
+      u = 16;
+    }
+    #pragma omp target exit data depend (inout: x, z) map (from: x, y, z, w)
+    if (x != 7 || y != 8 || z != 9 || s != 14 || t != 15 || u != 16 || w != 17)
+      abort ();
+  }
+  return 0;
+}

	Jakub