Hi Thomas!
On 07/09/2015 03:51 AM, Thomas Schwinge wrote:
Hi Jim!
On Tue, 7 Jul 2015 10:19:39 -0500, James Norris <jnor...@codesourcery.com>
wrote:
This patch fixes an issue where the deviceptr clause in an outer
directive was being ignored during implicit variable definition
on a nested directive.
Committed to gomp-4_0-branch.
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
I'm sorry, have not yet tried very hard; but I can't claim to understand
the logic here -- why is the OpenACC deviceptr clause so special? :-|
Because no data movement should be initiated on behalf of the user.
Prior to the fix, a map (tofrom) was being inserted when the variable
in question was implied within a nested directive. This is wrong.
No movement should occur as the memory for the variable is already
present on the target. Therefore, the use of 'present' is the correct
specification.
@@ -116,6 +116,9 @@ enum gimplify_omp_var_data
/* Gang-local OpenACC variable. */
GOVD_GANGLOCAL = (1 << 16),
+ /* OpenACC deviceptr clause. */
+ GOVD_USE_DEVPTR = (1 << 17),
+
GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
| GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
| GOVD_LOCAL)
@@ -6274,7 +6277,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
}
break;
}
+
flags = GOVD_MAP | GOVD_EXPLICIT;
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_DEVICEPTR)
+ flags |= GOVD_USE_DEVPTR;
goto do_add;
case OMP_CLAUSE_DEPEND:
@@ -6662,6 +6668,7 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void
*data)
: (flags & GOVD_FORCE_MAP
? GOMP_MAP_FORCE_TOFROM
: GOMP_MAP_TOFROM));
+
if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
{
@@ -6687,7 +6694,17 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void
*data)
OMP_CLAUSE_CHAIN (clause) = nc;
}
else
- OMP_CLAUSE_SIZE (clause) = DECL_SIZE_UNIT (decl);
+ {
+ if (gimplify_omp_ctxp->outer_context)
+ {
+ struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp->outer_context;
+ splay_tree_node on
+ = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
+ if (on && (on->value & GOVD_USE_DEVPTR))
+ OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_FORCE_PRESENT);
+ }
+ OMP_CLAUSE_SIZE (clause) = DECL_SIZE_UNIT (decl);
+ }
}
if (code == OMP_CLAUSE_FIRSTPRIVATE && (flags & GOVD_LASTPRIVATE) != 0)
{
The patch that you committed (r225518) also includes a test case
(thanks!) as follows:
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/deviceptr-4.c
@@ -0,0 +1,12 @@
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+void
+subr (int *a)
+{
+#pragma acc data deviceptr (a)
+#pragma acc parallel
+ a[0] += 1.0;
+}
+
+/* { dg-final { scan-tree-dump-times "#pragma omp target
oacc_parallel.*map\\(force_present:a \\\[len: 8\\\]\\)" 1 "gimple" } } */
+/* { dg-final { cleanup-tree-dump "gimple" } } */
That len: 8 is obviously valid only for 64-bit configurations, so will
cause a FAIL on anything else.
Fixed and committed to gomp-4_0-branch.
Thank you, thank you,
Jim
diff --git a/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c b/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c
index 4f6184c..0fef364 100644
--- a/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c
+++ b/gcc/testsuite/c-c++-common/goacc/deviceptr-4.c
@@ -8,5 +8,5 @@ subr (int *a)
a[0] += 1.0;
}
-/* { dg-final { scan-tree-dump-times "#pragma omp target oacc_parallel.*map\\(force_present:a \\\[len: 8\\\]\\)" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp target oacc_parallel.*map\\(force_present:a" 1 "gimple" } } */
/* { dg-final { cleanup-tree-dump "gimple" } } */
