mirror/gcc
2
0
Fork 0
mirror of git://gcc.gnu.org/git/gcc.git synced 2025-03-22 15:00:55 +08:00
Code Issues Packages Projects Releases Wiki Activity

Further changes for the OpenACC 'if_present' clause on the 'host_data' construct

Browse Source 
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.

From-SVN: r280149
This commit is contained in:
Thomas Schwinge 2020-01-10 23:23:44 +01:00 committed by Thomas Schwinge
parent 68be73fc42
commit b3b75e664a
11 changed files with 133 additions and 55 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

11
gcc/ChangeLog
View File

@ -1,3 +1,14 @@
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.

15
gcc/gimplify.c
View File

@ -12802,12 +12802,19 @@ 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,
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_HOST_DATA:
stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_HOST_DATA,
case OACC_KERNELS:
stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_KERNELS,
OMP_CLAUSES (expr));
break;
case OACC_PARALLEL:

12
gcc/omp-low.c
View File

@ -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);

5
gcc/testsuite/ChangeLog
View File

@ -1,3 +1,8 @@
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

38
gcc/testsuite/c-c++-common/goacc/host_data-1.c
View File

@ -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);
}
}
}

16
gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95
View File

@ -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" } }

3
gcc/tree-core.h
View File

@ -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

16
gcc/tree-pretty-print.c
View File

@ -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);

5
gcc/tree.h
View File

@ -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)

6
libgomp/ChangeLog
View File

@ -1,3 +1,9 @@
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

61
libgomp/target.c
View File

@ -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: