diff mbox

[gomp4.1] Depend clause support for offloading

Message ID 20150903141650.GX1847@tucnak.redhat.com
State New
Headers show

Commit Message

Jakub Jelinek Sept. 3, 2015, 2:16 p.m. UTC 
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.


	Jakub

Comments

Jakub Jelinek Sept. 3, 2015, 5:33 p.m. UTC | #1 
Hi!

FYI, I've merged trunk into the gomp-4_1-branch, it has been a while since
that has been done.  make -C check RUNTESTFLAGS=gomp.exp and
make check-target-libgomp still pass without offloading and when offloading
to mic emul (the latter with the libgomp.c/for-5.c and libgomp.c++/for-13.C
LTO ICEs that have been failing for a while).

	Jakub
diff mbox

Patch

--- 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;
+}