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
[OpenACC] Initial Manual Deep Copy 2018-10-02 Cesar Philippidis <ce...@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; +}