diff mbox series

[OpenACC] initial manual deep copy in c

Message ID 4d02c5f8-7b44-9d86-f78f-0925251a8418@codesourcery.com
State New
Headers show
Series [OpenACC] initial manual deep copy in c | expand

Commit Message

Cesar Philippidis Oct. 2, 2018, 10:14 p.m. UTC 
I've push the attach patch to my github trunk-acc-mdc branch which
enables OpenMP 4.5 deep copy semantics in OpenACC data clauses in C. Now
GCC accepts data clauses of the form

  #pragma acc data copy(v.a[:n], v.b)

I think there are a couple of limitations in OpenMP that's going to
force me to introduce a new GOMP_MAP_ACC_STRUCT map kind. Basically,
GOMP_MAP_STRUCT reserves the minimum amount of device storage to the
member actually used in a struct. OpenACC allows the users to
dynamically attach and detach struct members, so GOMP_MAP_ACC_STRUCT
would need reserve enough memory for the entire struct. This is also
necessary for cases like this

  struct {
    int *a, b, *c;
  } v;

  #pragma acc data copy(v.b)
  {
    #pragma acc parallel copy(v.a[:n], v.c[:n])
  }

If the acc data directive is replaced with omp target data, and the acc
parallel replaced with omp target something, then the runtime would
crash because struct v has been partially mapped already.

Going forward, OpenACC 2.6 requires the runtime to maintain an
attachment counter to keep track if struct fields have been mapped. So
that's another justification for the GOMP_MAP_ACC_STRUCT type.

This is all an early work in progress. I'm still experimenting with some
other functionality. If you checkout that branch, beware it may be rebased.

Cesar
diff mbox series

Patch

[OpenACC] Initial Manual Deep Copy

2018-10-02  Cesar Philippidis  <cesar@codesourcery.com>

	gcc/c/
	* c-typeck.c (handle_omp_array_sections_1): Enable structs in acc
	data clauses.
	(c_finish_omp_clauses): Likewise.

	libgomp/
	* libgomp.h: Declare gomp_map_val.
	* oacc-parallel.c (GOACC_parallel_keyed): Use it to set devaddrs.
	* target.c (gomp_map_val): Remove static inline.
	* testsuite/libgomp.oacc-c-c++-common/deep-copy-1.c: New test.


diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index 9d09b8d65fd..0428f48952a 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -12605,7 +12605,6 @@  handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 	  return error_mark_node;
 	}
       if (TREE_CODE (t) == COMPONENT_REF
-	  && ort == C_ORT_OMP
 	  && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 	      || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO
 	      || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM))
@@ -13799,7 +13798,6 @@  c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	      break;
 	    }
 	  if (TREE_CODE (t) == COMPONENT_REF
-	      && (ort & C_ORT_OMP)
 	      && OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_)
 	    {
 	      if (DECL_BIT_FIELD (TREE_OPERAND (t, 1)))
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index 3a8cc2bd7d6..553d1bb81ba 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -996,6 +996,7 @@  extern void gomp_acc_insert_pointer (size_t, void **, size_t *, void *);
 extern void gomp_acc_remove_pointer (void *, size_t, bool, int, int, int);
 extern void gomp_acc_declare_allocate (bool, size_t, void **, size_t *,
 				       unsigned short *);
+extern uintptr_t gomp_map_val (struct target_mem_desc *, void **, size_t);
 
 extern struct target_mem_desc *gomp_map_vars (struct gomp_device_descr *,
 					      size_t, void **, void **,
diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c
index b80ace58590..fd5bbfbdf7d 100644
--- a/libgomp/oacc-parallel.c
+++ b/libgomp/oacc-parallel.c
@@ -231,8 +231,7 @@  GOACC_parallel_keyed (int device, void (*fn) (void *),
 
   devaddrs = gomp_alloca (sizeof (void *) * mapnum);
   for (i = 0; i < mapnum; i++)
-    devaddrs[i] = (void *) (tgt->list[i].key->tgt->tgt_start
-			    + tgt->list[i].key->tgt_offset);
+    devaddrs[i] = (void *) gomp_map_val (tgt, hostaddrs, i);
 
   acc_dev->openacc.exec_func (tgt_fn, mapnum, hostaddrs, devaddrs,
 			      async, dims, tgt);
diff --git a/libgomp/target.c b/libgomp/target.c
index dda041cdbef..a87ba7cad0e 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -457,7 +457,7 @@  gomp_map_fields_existing (struct target_mem_desc *tgt, splay_tree_key n,
 	      (void *) cur_node.host_end);
 }
 
-static inline uintptr_t
+uintptr_t
 gomp_map_val (struct target_mem_desc *tgt, void **hostaddrs, size_t i)
 {
   if (tgt->list[i].key != NULL)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-1.c
new file mode 100644
index 00000000000..d489cc645cd
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-1.c
@@ -0,0 +1,25 @@ 
+#include <stdio.h>
+#include <stdlib.h>
+
+struct dc
+{
+  int a;
+  int *b;
+};
+
+int
+main ()
+{
+  int n = 100, i;
+  struct dc v = { .a = 3, .b = (int *) malloc (sizeof (int) * n) };
+
+#pragma omp target teams distribute parallel for map(tofrom:v.a, v.b[:n])
+#pragma acc parallel loop copy(v.a, v.b[:n])
+  for (i = 0; i < n; i++)
+    v.b[i] = v.a;
+
+  for (i = 0; i < 10; i++)
+    printf ("%d: %d\n", i, v.b[i]);
+
+  return 0;
+}