+2020-01-10 Thomas Schwinge <thomas@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 <dmalcolm@redhat.com>
* Makefile.in (OBJS): Add tree-diagnostic-path.o.
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));
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);
+2020-01-10 Thomas Schwinge <thomas@codesourcery.com>
+
+ * c-c++-common/goacc/host_data-1.c: Extend.
+ * gfortran.dg/goacc/host_data-tree.f95: Likewise.
+
2020-01-10 Jakub Jelinek <jakub@redhat.com>
PR tree-optimization/93210
/* 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" } } */
;
}
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])
/* 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);
}
}
}
-! { 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" } }
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
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:
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;
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";
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);
#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)
+2020-01-10 Thomas Schwinge <thomas@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 <jakub@redhat.com>
PR libgomp/93219
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;
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:
git://gcc.gnu.org / gcc.git / commitdiff