* [gomp4] Handle deviceptr from an outer directive
@ 2015-07-07 15:19 James Norris
2015-07-09 8:51 ` Thomas Schwinge
0 siblings, 1 reply; 3+ messages in thread
From: James Norris @ 2015-07-07 15:19 UTC (permalink / raw)
To: gcc-patches; +Cc: Thomas Schwinge
[-- Attachment #1: Type: text/plain, Size: 193 bytes --]
Hi,
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.
Jim
[-- Attachment #2: deviceptr.patch --]
[-- Type: text/x-patch, Size: 1697 bytes --]
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 51aadc0..a721a52 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -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)
{
^ permalink raw reply [flat|nested] 3+ messages in thread
* Re: [gomp4] Handle deviceptr from an outer directive
2015-07-07 15:19 [gomp4] Handle deviceptr from an outer directive James Norris
@ 2015-07-09 8:51 ` Thomas Schwinge
2015-07-10 17:40 ` James Norris
0 siblings, 1 reply; 3+ messages in thread
From: Thomas Schwinge @ 2015-07-09 8:51 UTC (permalink / raw)
To: James Norris; +Cc: gcc-patches, Jakub Jelinek
[-- Attachment #1: Type: text/plain, Size: 2931 bytes --]
Hi Jim!
On Tue, 7 Jul 2015 10:19:39 -0500, James Norris <jnorris@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? :-|
> @@ -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.
Grüße,
Thomas
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]
^ permalink raw reply [flat|nested] 3+ messages in thread
* Re: [gomp4] Handle deviceptr from an outer directive
2015-07-09 8:51 ` Thomas Schwinge
@ 2015-07-10 17:40 ` James Norris
0 siblings, 0 replies; 3+ messages in thread
From: James Norris @ 2015-07-10 17:40 UTC (permalink / raw)
To: Thomas Schwinge; +Cc: gcc-patches, Jakub Jelinek
[-- Attachment #1: Type: text/plain, Size: 3457 bytes --]
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 <jnorris@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
[-- Attachment #2: test.patch --]
[-- Type: text/x-patch, Size: 593 bytes --]
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" } } */
^ permalink raw reply [flat|nested] 3+ messages in thread
end of thread, other threads:[~2015-07-10 17:40 UTC | newest]
Thread overview: 3+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2015-07-07 15:19 [gomp4] Handle deviceptr from an outer directive James Norris
2015-07-09 8:51 ` Thomas Schwinge
2015-07-10 17:40 ` James Norris
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).