Hi! On 2020-01-09T22:00:09+0100, I wrote: > On 2019-12-24T15:23:56+0100, Tobias Burnus <tob...@codesourcery.com> wrote: >> PS: History: [...] >> A minor fix was committed to OG8?/OG9 as Rev. >> 995f9680a46c3a7246fe465faa847f8009e47ed8. > > That patch by Julian is on og9 only, and -- as far as I can tell -- has > never been posted/discussed. It may be obvious enough; we shall look > into it later. (I have not yet verified whether it is sufficiently > contained in the patch you posted here.)
Done.
>> --- a/gcc/omp-low.c
>> +++ b/gcc/omp-low.c
>> @@ -12006,6 +12006,9 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
>> omp_context *ctx)
> | case OMP_CLAUSE_USE_DEVICE_PTR:
> | case OMP_CLAUSE_USE_DEVICE_ADDR:
> | case OMP_CLAUSE_IS_DEVICE_PTR:
> | ovar = OMP_CLAUSE_DECL (c);
> | var = lookup_decl_in_outer_ctx (ovar, ctx);
> |
> | if (lang_hooks.decls.omp_array_data (ovar, true))
> | {
> | tkind = (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR
> | ? GOMP_MAP_USE_DEVICE_PTR : GOMP_MAP_FIRSTPRIVATE_INT);
> | x = build_sender_ref ((splay_tree_key) &DECL_NAME (ovar), ctx);
> | }
> | else if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR)
> | {
> | tkind = GOMP_MAP_USE_DEVICE_PTR;
> | x = build_sender_ref ((splay_tree_key) &DECL_UID (ovar), ctx);
> | }
> | else
> | {
>> tkind = GOMP_MAP_FIRSTPRIVATE_INT;
>> x = build_sender_ref (ovar, ctx);
>> }
>> + if (tkind == GOMP_MAP_USE_DEVICE_PTR
>> + && omp_find_clause (clauses, OMP_CLAUSE_IF_PRESENT))
>> + tkind = GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT;
>
> (For context: OpenACC only has an 'use_device' clause, mapped to
> 'OMP_CLAUSE_USE_DEVICE_PTR'; the other ones are for OpenMP only.)
> Can you foresee any C/C++/Fortran "magic" ;-) such that for OpenACC
> 'host_data' construct with 'use_device' clause(s), we'd get something
> different from 'tkind == GOMP_MAP_USE_DEVICE_PTR' here?
Tobias is looking into that.
> In that case,
> any 'if_present' clause would not be effective. (I have a WIP patch to
> 'gcc_assert' that.)
Done; doesn't trigger given the current test cases.
Also, calling 'omp_find_clause' here in a loop (over all mappings) is
something that Jakub would normally flag up ;-) -- I reworked this code
to use the standard idiom of doing such a transformation (here:
'if_present' clause applies to all 'use_device' clauses) early, in the
standard place in the gimplifier, by means of a new
'OMP_CLAUSE_USE_DEVICE_PTR_IF_PRESENT' flag, and also add more testsuite
coverage.
>> --- a/libgomp/target.c
>> +++ b/libgomp/target.c
>
> I'm have a WIP patch that should make this even more simpler.
Done. This is mostly code restructuring, but also, in the first loop,
"not present" case, we set 'tgt->list[i].offset = OFFSET_INLINED'
already, to avoid again processing this in the second loop.
> If that makes sense to consider (now or later), as I see there's some
> overlap in the 'gcc/omp-low.c:lower_omp_target' code paths: do we need
> test cases to verify 'if_present' in combination with Fortran optional
> arguments?
Tobias is looking into that.
See attached "Further changes for the OpenACC 'if_present' clause on the
'host_data' construct"; committed to trunk in r280149.
Grüße
Thomas
From 36c73c7d502ead3b941570a12a45e2cbe94ef53f Mon Sep 17 00:00:00 2001
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Fri, 10 Jan 2020 22:23:44 +0000
Subject: [PATCH] Further changes for the OpenACC 'if_present' clause on the
'host_data' construct
gcc/
* tree.h (OMP_CLAUSE_USE_DEVICE_PTR_IF_PRESENT): New definition.
* tree-core.h: Document it.
* gimplify.c (gimplify_omp_workshare): Set it.
* omp-low.c (lower_omp_target): Use it.
* tree-pretty-print.c (dump_omp_clause): Print it.
gcc/testsuite/
* c-c++-common/goacc/host_data-1.c: Extend.
* gfortran.dg/goacc/host_data-tree.f95: Likewise.
gcc/
* omp-low.c (lower_omp_target) <OMP_CLAUSE_USE_DEVICE_PTR etc.>:
Assert that for OpenACC we always have 'GOMP_MAP_USE_DEVICE_PTR'.
libgomp/
* target.c (gomp_map_vars_internal)
<GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT>: Clean up/elaborate code
paths.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@280149 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog | 11 ++++
gcc/gimplify.c | 15 +++--
gcc/omp-low.c | 12 +++-
gcc/testsuite/ChangeLog | 5 ++
.../c-c++-common/goacc/host_data-1.c | 38 +++++++++---
.../gfortran.dg/goacc/host_data-tree.f95 | 16 ++---
gcc/tree-core.h | 3 +
gcc/tree-pretty-print.c | 16 ++---
gcc/tree.h | 5 ++
libgomp/ChangeLog | 6 ++
libgomp/target.c | 61 ++++++++++++-------
11 files changed, 133 insertions(+), 55 deletions(-)
diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 9e9f8221af4..a195863212e 100644
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,14 @@
+2020-01-10 Thomas Schwinge <tho...@codesourcery.com>
+
+ * tree.h (OMP_CLAUSE_USE_DEVICE_PTR_IF_PRESENT): New definition.
+ * tree-core.h: Document it.
+ * gimplify.c (gimplify_omp_workshare): Set it.
+ * omp-low.c (lower_omp_target): Use it.
+ * tree-pretty-print.c (dump_omp_clause): Print it.
+
+ * omp-low.c (lower_omp_target) <OMP_CLAUSE_USE_DEVICE_PTR etc.>:
+ Assert that for OpenACC we always have 'GOMP_MAP_USE_DEVICE_PTR'.
+
2020-01-10 David Malcolm <dmalc...@redhat.com>
* Makefile.in (OBJS): Add tree-diagnostic-path.o.
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 00d264fc90f..fe7236de4c3 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -12802,14 +12802,21 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_DATA,
OMP_CLAUSES (expr));
break;
- case OACC_KERNELS:
- stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_KERNELS,
- OMP_CLAUSES (expr));
- break;
case OACC_HOST_DATA:
+ if (omp_find_clause (OMP_CLAUSES (expr), OMP_CLAUSE_IF_PRESENT))
+ {
+ for (tree c = OMP_CLAUSES (expr); c; c = OMP_CLAUSE_CHAIN (c))
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR)
+ OMP_CLAUSE_USE_DEVICE_PTR_IF_PRESENT (c) = 1;
+ }
+
stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_HOST_DATA,
OMP_CLAUSES (expr));
break;
+ case OACC_KERNELS:
+ stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_KERNELS,
+ OMP_CLAUSES (expr));
+ break;
case OACC_PARALLEL:
stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_PARALLEL,
OMP_CLAUSES (expr));
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 9a36192e8ef..eb3fe9688fe 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -12006,9 +12006,15 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
tkind = GOMP_MAP_FIRSTPRIVATE_INT;
x = build_sender_ref (ovar, ctx);
}
- if (tkind == GOMP_MAP_USE_DEVICE_PTR
- && omp_find_clause (clauses, OMP_CLAUSE_IF_PRESENT))
- tkind = GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT;
+
+ if (is_gimple_omp_oacc (ctx->stmt))
+ {
+ gcc_assert (tkind == GOMP_MAP_USE_DEVICE_PTR);
+
+ if (OMP_CLAUSE_USE_DEVICE_PTR_IF_PRESENT (c))
+ tkind = GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT;
+ }
+
type = TREE_TYPE (ovar);
if (lang_hooks.decls.omp_array_data (ovar, true))
var = lang_hooks.decls.omp_array_data (ovar, false);
diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index 6576aee81c8..cccc2853ed5 100644
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,3 +1,8 @@
+2020-01-10 Thomas Schwinge <tho...@codesourcery.com>
+
+ * c-c++-common/goacc/host_data-1.c: Extend.
+ * gfortran.dg/goacc/host_data-tree.f95: Likewise.
+
2020-01-10 Jakub Jelinek <ja...@redhat.com>
PR tree-optimization/93210
diff --git a/gcc/testsuite/c-c++-common/goacc/host_data-1.c b/gcc/testsuite/c-c++-common/goacc/host_data-1.c
index 658b7a677bc..ac244461467 100644
--- a/gcc/testsuite/c-c++-common/goacc/host_data-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/host_data-1.c
@@ -1,14 +1,20 @@
/* Test valid use of host_data directive. */
+/* { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" } */
+
int v1[3][3];
void
f (void)
{
#pragma acc host_data use_device(v1)
+ /* { dg-final { scan-tree-dump-times "(?n)#pragma acc host_data use_device_ptr\\(v1\\)$" 1 "original" } }
+ { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_host_data use_device_ptr\\(v1\\)$" 1 "gimple" } } */
;
#pragma acc host_data use_device(v1) if_present
+ /* { dg-final { scan-tree-dump-times "(?n)#pragma acc host_data if_present use_device_ptr\\(v1\\)$" 1 "original" } }
+ { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_host_data if_present use_device_ptr\\(if_present:v1\\)$" 1 "gimple" } } */
;
}
@@ -16,7 +22,7 @@ f (void)
void bar (float *, float *);
void
-foo (float *x, float *y)
+foo (float *x, float *y, float *yy)
{
int n = 1 << 10;
#pragma acc data create(x[0:n])
@@ -25,26 +31,38 @@ foo (float *x, float *y)
/* This should fail at run time because y is not mapped. */
#pragma acc host_data use_device(x,y)
+ /* { dg-final { scan-tree-dump-times "(?n)#pragma acc host_data use_device_ptr\\(y\\) use_device_ptr\\(x\\)$" 1 "original" } }
+ { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_host_data use_device_ptr\\(y\\) use_device_ptr\\(x\\)$" 1 "gimple" } } */
bar (x, y);
/* y is still not mapped, but this should not fail at run time but
continue execution with y remaining as the host address. */
#pragma acc host_data use_device(x,y) if_present
+ /* { dg-final { scan-tree-dump-times "(?n)#pragma acc host_data if_present use_device_ptr\\(y\\) use_device_ptr\\(x\\)$" 1 "original" } }
+ { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_host_data if_present use_device_ptr\\(if_present:y\\) use_device_ptr\\(if_present:x\\)$" 1 "gimple" } } */
bar (x, y);
-#pragma acc data copyout(y[0:n])
+#pragma acc data copyout(yy[0:n])
{
-#pragma acc host_data use_device(x,y)
- bar (x, y);
+#pragma acc host_data use_device(x,yy)
+ /* { dg-final { scan-tree-dump-times "(?n)#pragma acc host_data use_device_ptr\\(yy\\) use_device_ptr\\(x\\)$" 1 "original" } }
+ { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_host_data use_device_ptr\\(yy\\) use_device_ptr\\(x\\)$" 1 "gimple" } } */
+ bar (x, yy);
-#pragma acc host_data use_device(x,y) if_present
- bar (x, y);
+#pragma acc host_data use_device(x,yy) if_present
+ /* { dg-final { scan-tree-dump-times "(?n)#pragma acc host_data if_present use_device_ptr\\(yy\\) use_device_ptr\\(x\\)$" 1 "original" } }
+ { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_host_data if_present use_device_ptr\\(if_present:yy\\) use_device_ptr\\(if_present:x\\)$" 1 "gimple" } } */
+ bar (x, yy);
-#pragma acc host_data use_device(x,y) if(x != y)
- bar (x, y);
+#pragma acc host_data use_device(x,yy) if(x != yy)
+ /* { dg-final { scan-tree-dump-times "(?n)#pragma acc host_data if\\(x \\!= yy\\) use_device_ptr\\(yy\\) use_device_ptr\\(x\\)$" 1 "original" } }
+ { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_host_data if\\(D\\.\[0-9\]+\\) use_device_ptr\\(yy\\) use_device_ptr\\(x\\)$" 1 "gimple" } } */
+ bar (x, yy);
-#pragma acc host_data use_device(x,y) if_present if(x != y)
- bar (x, y);
+#pragma acc host_data use_device(x,yy) if_present if(x == yy)
+ /* { dg-final { scan-tree-dump-times "(?n)#pragma acc host_data if\\(x == yy\\) if_present use_device_ptr\\(yy\\) use_device_ptr\\(x\\)$" 1 "original" } }
+ { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_host_data if\\(D\\.\[0-9\]+\\) if_present use_device_ptr\\(if_present:yy\\) use_device_ptr\\(if_present:x\\)$" 1 "gimple" } } */
+ bar (x, yy);
}
}
}
diff --git a/gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95
index 2ac1c0d66d6..558e80014d7 100644
--- a/gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95
@@ -1,21 +1,23 @@
-! { dg-do compile }
-! { dg-additional-options "-fdump-tree-original" }
+! { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" }
program test
implicit none
integer, pointer :: p
!$acc host_data use_device(p)
+ ! { dg-final { scan-tree-dump-times "(?n)#pragma acc host_data use_device_ptr\\(p\\)$" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_host_data use_device_ptr\\(p\\)$" 1 "gimple" } }
!$acc end host_data
!$acc host_data use_device(p) if (p == 42)
+ ! { dg-final { scan-tree-dump-times "(?n)D\\.\[0-9\]+ = \\*p == 42;$" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "(?n)#pragma acc host_data use_device_ptr\\(p\\) if\\(D\\.\[0-9\]+\\)$" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_host_data use_device_ptr\\(p\\) if\\(D\\.\[0-9\]+\\)$" 1 "gimple" } }
!$acc end host_data
!$acc host_data use_device(p) if_present if (p == 43)
+ ! { dg-final { scan-tree-dump-times "(?n)D\\.\[0-9\]+ = \\*p == 43;$" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "(?n)#pragma acc host_data use_device_ptr\\(p\\) if\\(D\\.\[0-9\]+\\) if_present$" 1 "original" } }
+ ! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_host_data use_device_ptr\\(if_present:p\\) if\\(D\\.\[0-9\]+\\) if_present$" 1 "gimple" } }
!$acc end host_data
end program test
-! { dg-final { scan-tree-dump-times "pragma acc host_data use_device_ptr\\(p\\)" 3 "original" } }
-! { dg-final { scan-tree-dump-times "D.\[0-9\]+ = \\*p == 42;" 1 "original" } }
-! { dg-final { scan-tree-dump-times "pragma acc host_data use_device_ptr\\(p\\) if\\(D.\[0-9\]+\\)" 2 "original" } }
-! { dg-final { scan-tree-dump-times "D.\[0-9\]+ = \\*p == 43;" 1 "original" } }
-! { dg-final { scan-tree-dump-times "pragma acc host_data use_device_ptr\\(p\\) if\\(D.\[0-9\]+\\) if_present" 1 "original" } }
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index 62130f488cc..765ea2a9542 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -1175,6 +1175,9 @@ struct GTY(()) tree_base {
OMP_CLAUSE_REDUCTION_OMP_ORIG_REF in
OMP_CLAUSE_{,TASK_,IN_}REDUCTION
+ OMP_CLAUSE_USE_DEVICE_PTR_IF_PRESENT in
+ OMP_CLAUSE_USE_DEVICE_PTR
+
TRANSACTION_EXPR_RELAXED in
TRANSACTION_EXPR
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index e895a4f6609..fe2e62b31ba 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -432,7 +432,7 @@ static void
dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
{
const char *name;
-
+ const char *modifier = NULL;
switch (OMP_CLAUSE_CODE (clause))
{
case OMP_CLAUSE_PRIVATE:
@@ -446,13 +446,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
goto print_remap;
case OMP_CLAUSE_LASTPRIVATE:
name = "lastprivate";
- if (!OMP_CLAUSE_LASTPRIVATE_CONDITIONAL (clause))
- goto print_remap;
- pp_string (pp, "lastprivate(conditional:");
- dump_generic_node (pp, OMP_CLAUSE_DECL (clause),
- spc, flags, false);
- pp_right_paren (pp);
- break;
+ if (OMP_CLAUSE_LASTPRIVATE_CONDITIONAL (clause))
+ modifier = "conditional:";
+ goto print_remap;
case OMP_CLAUSE_COPYIN:
name = "copyin";
goto print_remap;
@@ -464,6 +460,8 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
goto print_remap;
case OMP_CLAUSE_USE_DEVICE_PTR:
name = "use_device_ptr";
+ if (OMP_CLAUSE_USE_DEVICE_PTR_IF_PRESENT (clause))
+ modifier = "if_present:";
goto print_remap;
case OMP_CLAUSE_USE_DEVICE_ADDR:
name = "use_device_addr";
@@ -501,6 +499,8 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
print_remap:
pp_string (pp, name);
pp_left_paren (pp);
+ if (modifier)
+ pp_string (pp, modifier);
dump_generic_node (pp, OMP_CLAUSE_DECL (clause),
spc, flags, false);
pp_right_paren (pp);
diff --git a/gcc/tree.h b/gcc/tree.h
index 9ca9ab58ec0..93422206b63 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1627,6 +1627,11 @@ class auto_suppress_location_wrappers
#define OMP_CLAUSE_MAP_IN_REDUCTION(NODE) \
TREE_PRIVATE (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP))
+/* True on an OMP_CLAUSE_USE_DEVICE_PTR with an OpenACC 'if_present'
+ clause. */
+#define OMP_CLAUSE_USE_DEVICE_PTR_IF_PRESENT(NODE) \
+ (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_USE_DEVICE_PTR)->base.public_flag)
+
#define OMP_CLAUSE_PROC_BIND_KIND(NODE) \
(OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_PROC_BIND)->omp_clause.subcode.proc_bind_kind)
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 152e52c21c1..81d0c164a3a 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,9 @@
+2020-01-10 Thomas Schwinge <tho...@codesourcery.com>
+
+ * target.c (gomp_map_vars_internal)
+ <GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT>: Clean up/elaborate code
+ paths.
+
2020-01-10 Jakub Jelinek <ja...@redhat.com>
PR libgomp/93219
diff --git a/libgomp/target.c b/libgomp/target.c
index 522b69e6d5d..38de1c0cf92 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -740,22 +740,24 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
cur_node.host_start = (uintptr_t) hostaddrs[i];
cur_node.host_end = cur_node.host_start;
splay_tree_key n = gomp_map_lookup (mem_map, &cur_node);
- if (n == NULL)
+ if (n != NULL)
+ {
+ cur_node.host_start -= n->host_start;
+ hostaddrs[i]
+ = (void *) (n->tgt->tgt_start + n->tgt_offset
+ + cur_node.host_start);
+ }
+ else if ((kind & typemask) == GOMP_MAP_USE_DEVICE_PTR)
{
- if ((kind & typemask) == GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT)
- {
- /* If not present, continue using the host address. */
- tgt->list[i].offset = 0;
- continue;
- }
gomp_mutex_unlock (&devicep->lock);
gomp_fatal ("use_device_ptr pointer wasn't mapped");
}
- cur_node.host_start -= n->host_start;
- hostaddrs[i]
- = (void *) (n->tgt->tgt_start + n->tgt_offset
- + cur_node.host_start);
- tgt->list[i].offset = ~(uintptr_t) 0;
+ else if ((kind & typemask) == GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT)
+ /* If not present, continue using the host address. */
+ ;
+ else
+ __builtin_unreachable ();
+ tgt->list[i].offset = OFFSET_INLINED;
}
else
tgt->list[i].offset = 0;
@@ -980,27 +982,40 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
case GOMP_MAP_FIRSTPRIVATE_INT:
case GOMP_MAP_ZERO_LEN_ARRAY_SECTION:
continue;
- case GOMP_MAP_USE_DEVICE_PTR:
case GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT:
+ /* The OpenACC 'host_data' construct only allows 'use_device'
+ "mapping" clauses, so in the first loop, 'not_found_cnt'
+ must always have been zero, so all OpenACC 'use_device'
+ clauses have already been handled. (We can only easily test
+ 'use_device' with 'if_present' clause here.) */
+ assert (tgt->list[i].offset == OFFSET_INLINED);
+ /* Nevertheless, FALLTHRU to the normal handling, to keep the
+ code conceptually simple, similar to the first loop. */
+ case GOMP_MAP_USE_DEVICE_PTR:
if (tgt->list[i].offset == 0)
{
cur_node.host_start = (uintptr_t) hostaddrs[i];
cur_node.host_end = cur_node.host_start;
n = gomp_map_lookup (mem_map, &cur_node);
- if (n == NULL)
+ if (n != NULL)
+ {
+ cur_node.host_start -= n->host_start;
+ hostaddrs[i]
+ = (void *) (n->tgt->tgt_start + n->tgt_offset
+ + cur_node.host_start);
+ }
+ else if ((kind & typemask) == GOMP_MAP_USE_DEVICE_PTR)
{
- if ((kind & typemask)
- == GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT)
- /* If not present, continue using the host address. */
- continue;
gomp_mutex_unlock (&devicep->lock);
gomp_fatal ("use_device_ptr pointer wasn't mapped");
}
- cur_node.host_start -= n->host_start;
- hostaddrs[i]
- = (void *) (n->tgt->tgt_start + n->tgt_offset
- + cur_node.host_start);
- tgt->list[i].offset = ~(uintptr_t) 0;
+ else if ((kind & typemask)
+ == GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT)
+ /* If not present, continue using the host address. */
+ ;
+ else
+ __builtin_unreachable ();
+ tgt->list[i].offset = OFFSET_INLINED;
}
continue;
case GOMP_MAP_STRUCT:
--
2.17.1
signature.asc
Description: PGP signature
