2018-XX-YY Cesar Philippidis <cesar@codesourcery.com>
gcc/
* gimplify.c (gimplify_adjust_omp_clauses): Filter out
GOMP_MAP_STRUCT for acc exit data.
(gimplify_omp_target_update): Promote GOMP_MAP_DETACH
to GOMP_MAP_FORCE_DETACH when the finalize clause is present.
* omp-low.c (lower_omp_target): Add support for GOMP_MAP_{ATTACH,
DETACH, FORCE_DETACH}.
* tree-pretty-print.c (dump_omp_clause): Likewise.
gcc/c-family/
* c-pragma.h (enum pragma_omp_clause): Define
PRAGMA_OACC_CLAUSE_{ATTACH,DETACH}.
include/
* gomp-constants.h (GOMP_MAP_DEEP_COPY): Define.
(enum gomp_map_kind): Add GOMP_MAP_{ATTACH, DETACH, FORCE_DETACH}.
@@ -131,11 +131,13 @@ enum pragma_omp_clause {
/* Clauses for OpenACC. */
PRAGMA_OACC_CLAUSE_ASYNC,
+ PRAGMA_OACC_CLAUSE_ATTACH,
PRAGMA_OACC_CLAUSE_AUTO,
PRAGMA_OACC_CLAUSE_COPY,
PRAGMA_OACC_CLAUSE_COPYOUT,
PRAGMA_OACC_CLAUSE_CREATE,
PRAGMA_OACC_CLAUSE_DELETE,
+ PRAGMA_OACC_CLAUSE_DETACH,
PRAGMA_OACC_CLAUSE_DEVICEPTR,
PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT,
PRAGMA_OACC_CLAUSE_FINALIZE,
@@ -9145,7 +9145,8 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
}
}
else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT
- && code == OMP_TARGET_EXIT_DATA)
+ && (code == OMP_TARGET_EXIT_DATA
+ || code == OACC_EXIT_DATA))
remove = true;
else if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST
@@ -11001,8 +11002,9 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
&& omp_find_clause (OMP_STANDALONE_CLAUSES (expr),
OMP_CLAUSE_FINALIZE))
{
- /* Use GOMP_MAP_DELETE/GOMP_MAP_FORCE_FROM to denote that "finalize"
- semantics apply to all mappings of this OpenACC directive. */
+ /* Use GOMP_MAP_DELETE, GOMP_MAP_FORCE_DETACH, and
+ GOMP_MAP_FORCE_FROM to denote that "finalize" semantics apply
+ to all mappings of this OpenACC directive. */
bool finalize_marked = false;
for (tree c = OMP_STANDALONE_CLAUSES (expr); c; c = OMP_CLAUSE_CHAIN (c))
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
@@ -11016,6 +11018,10 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_DELETE);
finalize_marked = true;
break;
+ case GOMP_MAP_DETACH:
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_DETACH);
+ finalize_marked = true;
+ break;
default:
/* Check consistency: libgomp relies on the very first data
mapping clause being marked, so make sure we did that before
@@ -7581,6 +7581,9 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
case GOMP_MAP_FORCE_DEVICEPTR:
case GOMP_MAP_DEVICE_RESIDENT:
case GOMP_MAP_LINK:
+ case GOMP_MAP_ATTACH:
+ case GOMP_MAP_DETACH:
+ case GOMP_MAP_FORCE_DETACH:
gcc_assert (is_gimple_omp_oacc (stmt));
break;
default:
@@ -755,6 +755,15 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
case GOMP_MAP_LINK:
pp_string (pp, "link");
break;
+ case GOMP_MAP_ATTACH:
+ pp_string (pp, "attach");
+ break;
+ case GOMP_MAP_DETACH:
+ pp_string (pp, "detach");
+ break;
+ case GOMP_MAP_FORCE_DETACH:
+ pp_string (pp, "force_detach");
+ break;
default:
gcc_unreachable ();
}
@@ -42,6 +42,7 @@
#define GOMP_MAP_FLAG_SPECIAL_2 (1 << 4)
#define GOMP_MAP_FLAG_SPECIAL (GOMP_MAP_FLAG_SPECIAL_1 \
| GOMP_MAP_FLAG_SPECIAL_0)
+#define GOMP_MAP_DEEP_COPY (1 << 5)
/* Flag to force a specific behavior (or else, trigger a run-time error). */
#define GOMP_MAP_FLAG_FORCE (1 << 7)
@@ -128,6 +129,13 @@ enum gomp_map_kind
/* Decrement usage count and deallocate if zero. */
GOMP_MAP_RELEASE = (GOMP_MAP_FLAG_SPECIAL_2
| GOMP_MAP_DELETE),
+ /* In OpenACC, attach a pointer to a mapped struct field. */
+ GOMP_MAP_ATTACH = (GOMP_MAP_DEEP_COPY | 0),
+ /* In OpenACC, detach a pointer to a mapped struct field. */
+ GOMP_MAP_DETACH = (GOMP_MAP_DEEP_COPY | 1),
+ /* In OpenACC, detach a pointer to a mapped struct field. */
+ GOMP_MAP_FORCE_DETACH = (GOMP_MAP_DEEP_COPY
+ | GOMP_MAP_FLAG_FORCE | 1),
/* Internal to GCC, not used in libgomp. */
/* Do not map, but pointer assign a pointer instead. */
--------------2.17.1--
This patch series adds support for the new attach / detach clauses introduced in OpenACC 2.6 to the C and C++ front ends. Julian is working patches for the Fortran front end along with the runtime. As their names somewhat imply, attach and detach are new data clauses that are used to support manual deep copy in OpenACC. Specifically, OpenACC 2.6 allows users to specify individual structure fields inside data clauses, whereas before that would only work inside the update directive. The attach and detach clauses allow users to update the pointers in structure fields with their on-device counterparts. As an example, consider the the following code: struct { int *a, b } s; int *z = ... #pragma acc enter data copyin(a[:N], s) ... s.a = z; #pragma acc enter data attach(s.a) ... #pragma acc exit data detach(s.a) #pragma acc exit data copyout(s) Because the attach clause updates field s.a with the device address, "acc exit data detach" must be used to restore the host pointer contents before that value is copied back to the host. This patch in particular adds the generic infrastructure for the attach and detach clauses. All of the front ends lower the attach clause as GOMP_MAP_DETACH data mapping. However, if a detachment is finalized, e.g. #pragma acc exit data finalize detach(ptr) the gimplifier will promote it to GOMP_MAP_FORCE_FINALIZE. Also, this patch teaches the gimplifier how to ignore GOMP_MAP_STRUCT for the target update constructs. Is this patch OK for trunk? I bootstrapped and regression tested it for x86_64 Linux with nvptx offloading. Thanks, Cesar