From: Thomas Schwinge Date: Fri, 10 Jan 2020 22:23:44 +0000 (+0100) Subject: Further changes for the OpenACC 'if_present' clause on the 'host_data' construct X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=b3b75e664a619dae98571a0b3ac8034f5fa7c2be;p=gcc.git 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) : Assert that for OpenACC we always have 'GOMP_MAP_USE_DEVICE_PTR'. libgomp/ * target.c (gomp_map_vars_internal) : Clean up/elaborate code paths. From-SVN: r280149 --- 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 + + * 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) : + Assert that for OpenACC we always have 'GOMP_MAP_USE_DEVICE_PTR'. + 2020-01-10 David Malcolm * 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 + + * c-c++-common/goacc/host_data-1.c: Extend. + * gfortran.dg/goacc/host_data-tree.f95: Likewise. + 2020-01-10 Jakub Jelinek 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 + + * target.c (gomp_map_vars_internal) + : Clean up/elaborate code + paths. + 2020-01-10 Jakub Jelinek 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: