I had previously posted this patch as part of a monster deviceptr patch
here <https://gcc.gnu.org/ml/gcc-patches/2018-06/msg01911.html>. This
patch breaks out the generic gimplifier changes. Essentially, with this
patch, the gimplifier will now transfer deviceptr data clauses using
GOMP_MAP_FORCE_DEVICEPTR.
Is this patch OK for trunk? It bootstrapped / regression tested cleanly
for x86_64 with nvptx offloading.
Thanks,
Cesar
>From b5cf37b795ce78c78f3f434ac6999f7094bd86aa Mon Sep 17 00:00:00 2001
From: Cesar Philippidis <ce...@codesourcery.com>
Date: Mon, 7 May 2018 08:23:48 -0700
Subject: [PATCH] [OpenACC] Update deviceptr handling
2018-XX-YY Cesar Philippidis <ce...@codesourcery.com>
gcc/fortran/
* trans-openmp.c (gfc_omp_finish_clause): Don't create pointer data
mappings for deviceptr clauses.
(gfc_trans_omp_clauses): Likewise.
gcc/
* gimplify.c (enum gimplify_omp_var_data): Add GOVD_DEVICETPR.
(omp_notice_variable): Add GOVD_DEVICEPTR attribute when appropriate.
(gimplify_scan_omp_clauses): Likewise.
(gimplify_adjust_omp_clauses_1): Set GOMP_MAP_FORCE_DEVICEPTR for
implicit deviceptr mappings.
gcc/testsuite/
* c-c++-common/goacc/deviceptr-4.c: Update expected data mapping.
(cherry picked from openacc-gcc-7-branch commit
d3de16b461545aac1925f0d7c2851c8c49a07d06 and commit
f0514fe1899666bb5b8ee52601f5d4263d4c4646)
---
gcc/fortran/trans-openmp.c | 9 +++++++++
gcc/gimplify.c | 12 +++++++++++-
gcc/testsuite/c-c++-common/goacc/deviceptr-4.c | 2 +-
3 files changed, 21 insertions(+), 2 deletions(-)
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index f038f4c..ca31c88 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -1060,6 +1060,8 @@ gfc_omp_finish_clause (tree c, gimple_seq *pre_p)
}
tree c2 = NULL_TREE, c3 = NULL_TREE, c4 = NULL_TREE;
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_DEVICEPTR)
+ return;
if (POINTER_TYPE_P (TREE_TYPE (decl)))
{
if (!gfc_omp_privatize_by_reference (decl)
@@ -2111,6 +2113,12 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
if (n->expr == NULL || n->expr->ref->u.ar.type == AR_FULL)
{
if (POINTER_TYPE_P (TREE_TYPE (decl))
+ && n->u.map_op == OMP_MAP_FORCE_DEVICEPTR)
+ {
+ OMP_CLAUSE_DECL (node) = decl;
+ goto finalize_map_clause;
+ }
+ else if (POINTER_TYPE_P (TREE_TYPE (decl))
&& (gfc_omp_privatize_by_reference (decl)
|| GFC_DECL_GET_SCALAR_POINTER (decl)
|| GFC_DECL_GET_SCALAR_ALLOCATABLE (decl)
@@ -2282,6 +2290,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
ptr2 = fold_convert (sizetype, ptr2);
OMP_CLAUSE_SIZE (node3)
= fold_build2 (MINUS_EXPR, sizetype, ptr, ptr2);
+ finalize_map_clause:;
}
switch (n->u.map_op)
{
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 4a109ae..bcf862f 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -105,6 +105,9 @@ enum gimplify_omp_var_data
/* Flag for GOVD_MAP: must be present already. */
GOVD_MAP_FORCE_PRESENT = 524288,
+ /* Flag for OpenACC deviceptrs. */
+ GOVD_DEVICEPTR = (1<<21),
+
GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
| GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
| GOVD_LOCAL)
@@ -7232,6 +7235,7 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
error ("variable %qE declared in enclosing "
"%<host_data%> region", DECL_NAME (decl));
nflags |= GOVD_MAP;
+ nflags |= (n2->value & GOVD_DEVICEPTR);
if (octx->region_type == ORT_ACC_DATA
&& (n2->value & GOVD_MAP_0LEN_ARRAY))
nflags |= GOVD_MAP_0LEN_ARRAY;
@@ -8213,6 +8217,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TO
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TOFROM)
flags |= GOVD_MAP_ALWAYS_TO;
+ else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_DEVICEPTR)
+ flags |= GOVD_DEVICEPTR;
goto do_add;
case OMP_CLAUSE_DEPEND:
@@ -8828,7 +8834,8 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
/* Not all combinations of these GOVD_MAP flags are actually valid. */
switch (flags & (GOVD_MAP_TO_ONLY
| GOVD_MAP_FORCE
- | GOVD_MAP_FORCE_PRESENT))
+ | GOVD_MAP_FORCE_PRESENT
+ | GOVD_DEVICEPTR))
{
case 0:
kind = GOMP_MAP_TOFROM;
@@ -8845,6 +8852,9 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
case GOVD_MAP_FORCE_PRESENT:
kind = GOMP_MAP_FORCE_PRESENT;
break;
+ case GOVD_DEVICEPTR:
+ kind = GOMP_MAP_FORCE_DEVICEPTR;
+ break;
default:
gcc_unreachable ();
}
diff --git a/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c b/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c
index db1b916..79a5162 100644
--- a/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c
+++ b/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c
@@ -8,4 +8,4 @@ subr (int *a)
a[0] += 1.0;
}
-/* { dg-final { scan-tree-dump-times "#pragma omp target oacc_parallel.*map\\(tofrom:a" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp target oacc_parallel.*map\\(force_deviceptr:a" 1 "gimple" } } */
--
2.7.4
