mirrors/gcc
1
0
Fork 0
mirror of https://gcc.gnu.org/git/gcc.git synced 2024-11-24 03:14:08 +08:00
Code Issues Actions Packages Projects Releases Wiki Activity

gimplify.c (omp_add_variable): Use GOVD_PRIVATE | GOVD_EXPLICIT for VLA helper variables on target data even if...

Browse Source 
* gimplify.c (omp_add_variable): Use GOVD_PRIVATE | GOVD_EXPLICIT
	for VLA helper variables on target data even if not GOVD_FIRSTPRIVATE.
	(gimplify_scan_omp_clauses): For OMP_CLAUSE_USE_DEVICE_* use just
	GOVD_EXPLICIT flags.
	(gimplify_omp_workshare): For OMP_TARGET_DATA move all
	OMP_CLAUSE_USE_DEVICE_* clauses to the end of clauses chain.
	* omp-low.c (scan_sharing_clauses): For OMP_CLAUSE_USE_DEVICE_*
	call install_var_field with mask 11 instead of 3.
	(lower_omp_target): For OMP_CLAUSE_USE_DEVICE_* use pass
	(splay_tree_key) &DECL_UID (var) to build_sender_ref instead of var.
gcc/c/
	* c-typeck.c (c_finish_omp_clauses): For C_ORT_OMP
	OMP_CLAUSE_USE_DEVICE_* clauses use oacc_reduction_head bitmap
	instead of generic_head to track duplicates.
gcc/cp/
	* semantics.c (finish_omp_clauses): For C_ORT_OMP
	OMP_CLAUSE_USE_DEVICE_* clauses use oacc_reduction_head bitmap
	instead of generic_head to track duplicates.
libgomp/
	* target.c (gomp_map_vars_internal): For GOMP_MAP_USE_DEVICE_PTR
	perform the lookup in the first loop only if !not_found_cnt, otherwise
	perform lookups for it in the second loop guarded with
	if (not_found_cnt || has_firstprivate).
	* testsuite/libgomp.c/target-37.c: New test.
	* testsuite/libgomp.c++/target-22.C: New test.

From-SVN: r274206
This commit is contained in:
Jakub Jelinek 2019-08-08 08:39:02 +02:00 committed by Jakub Jelinek
parent 34f3ec0757
commit 8860d2706d
11 changed files with 288 additions and 32 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

13
gcc/ChangeLog
View File

@ -1,3 +1,16 @@
2019-08-08 Jakub Jelinek <jakub@redhat.com>
* gimplify.c (omp_add_variable): Use GOVD_PRIVATE | GOVD_EXPLICIT
for VLA helper variables on target data even if not GOVD_FIRSTPRIVATE.
(gimplify_scan_omp_clauses): For OMP_CLAUSE_USE_DEVICE_* use just
GOVD_EXPLICIT flags.
(gimplify_omp_workshare): For OMP_TARGET_DATA move all
OMP_CLAUSE_USE_DEVICE_* clauses to the end of clauses chain.
* omp-low.c (scan_sharing_clauses): For OMP_CLAUSE_USE_DEVICE_*
call install_var_field with mask 11 instead of 3.
(lower_omp_target): For OMP_CLAUSE_USE_DEVICE_* use pass
(splay_tree_key) &DECL_UID (var) to build_sender_ref instead of var.
2019-08-07 Richard Sandiford <richard.sandiford@arm.com>
* config/aarch64/constraints.md (Z): Handle floating-point zeros too.

6
gcc/c/ChangeLog
View File

@ -1,3 +1,9 @@
2019-08-08 Jakub Jelinek <jakub@redhat.com>
* c-typeck.c (c_finish_omp_clauses): For C_ORT_OMP
OMP_CLAUSE_USE_DEVICE_* clauses use oacc_reduction_head bitmap
instead of generic_head to track duplicates.
2019-08-07 Jakub Jelinek <jakub@redhat.com>
* c-parser.c (c_parser_omp_clause_name): Parse use_device_addr clause.

15
gcc/c/c-typeck.c
View File

@ -13680,7 +13680,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
/* If ort == C_ORT_OMP_DECLARE_SIMD used as uniform_head instead. */
bitmap_initialize (&map_head, &bitmap_default_obstack);
bitmap_initialize (&map_field_head, &bitmap_default_obstack);
/* If ort == C_ORT_OMP used as nontemporal_head instead. */
/* If ort == C_ORT_OMP used as nontemporal_head or use_device_xxx_head
instead. */
bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack);
if (ort & C_ORT_ACC)
@ -14072,13 +14073,19 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
else if (ort == C_ORT_ACC
&& OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
else if ((ort == C_ORT_ACC
&& OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
|| (ort == C_ORT_OMP
&& (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR
|| (OMP_CLAUSE_CODE (c)
== OMP_CLAUSE_USE_DEVICE_ADDR))))
{
if (bitmap_bit_p (&oacc_reduction_head, DECL_UID (t)))
{
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in reduction clauses",
ort == C_ORT_ACC
? "%qD appears more than once in reduction clauses"
: "%qD appears more than once in data clauses",
t);
remove = true;
}

6
gcc/cp/ChangeLog
View File

@ -1,3 +1,9 @@
2019-08-08 Jakub Jelinek <jakub@redhat.com>
* semantics.c (finish_omp_clauses): For C_ORT_OMP
OMP_CLAUSE_USE_DEVICE_* clauses use oacc_reduction_head bitmap
instead of generic_head to track duplicates.
2019-08-07 Marek Polacek <polacek@redhat.com>
PR c++/81429 - wrong parsing of constructor with C++11 attribute.

15
gcc/cp/semantics.c
View File

@ -6148,7 +6148,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
/* If ort == C_ORT_OMP_DECLARE_SIMD used as uniform_head instead. */
bitmap_initialize (&map_head, &bitmap_default_obstack);
bitmap_initialize (&map_field_head, &bitmap_default_obstack);
/* If ort == C_ORT_OMP used as nontemporal_head instead. */
/* If ort == C_ORT_OMP used as nontemporal_head or use_device_xxx_head
instead. */
bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack);
if (ort & C_ORT_ACC)
@ -6406,13 +6407,19 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
else if (ort == C_ORT_ACC
&& OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
else if ((ort == C_ORT_ACC
&& OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
|| (ort == C_ORT_OMP
&& (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR
|| (OMP_CLAUSE_CODE (c)
== OMP_CLAUSE_USE_DEVICE_ADDR))))
{
if (bitmap_bit_p (&oacc_reduction_head, DECL_UID (t)))
{
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in reduction clauses",
ort == C_ORT_ACC
? "%qD appears more than once in reduction clauses"
: "%qD appears more than once in data clauses",
t);
remove = true;
}

32
gcc/gimplify.c
View File

@ -6932,8 +6932,10 @@ omp_add_variable (struct gimplify_omp_ctx *ctx, tree decl, unsigned int flags)
nflags = GOVD_MAP | GOVD_MAP_TO_ONLY | GOVD_EXPLICIT;
else if (flags & GOVD_PRIVATE)
nflags = GOVD_PRIVATE;
else if ((ctx->region_type & (ORT_TARGET | ORT_TARGET_DATA)) != 0
&& (flags & GOVD_FIRSTPRIVATE))
else if (((ctx->region_type & (ORT_TARGET | ORT_TARGET_DATA)) != 0
&& (flags & GOVD_FIRSTPRIVATE))
|| (ctx->region_type == ORT_TARGET_DATA
&& (flags & GOVD_DATA_SHARE_CLASS) == 0))
nflags = GOVD_PRIVATE | GOVD_EXPLICIT;
else
nflags = GOVD_FIRSTPRIVATE;
@ -9016,6 +9018,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_USE_DEVICE_ADDR:
flags = GOVD_EXPLICIT;
goto do_add;
case OMP_CLAUSE_IS_DEVICE_PTR:
flags = GOVD_FIRSTPRIVATE | GOVD_EXPLICIT;
goto do_add;
@ -12404,8 +12409,27 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
OMP_CLAUSES (expr));
break;
case OMP_TARGET_DATA:
stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_DATA,
OMP_CLAUSES (expr));
/* Put use_device_{ptr,addr} clauses last, as map clauses are supposed
to be evaluated before the use_device_{ptr,addr} clauses if they
refer to the same variables. */
{
tree use_device_clauses;
tree *pc, *uc = &use_device_clauses;
for (pc = &OMP_CLAUSES (expr); *pc; )
if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_USE_DEVICE_PTR
|| OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_USE_DEVICE_ADDR)
{
*uc = *pc;
*pc = OMP_CLAUSE_CHAIN (*pc);
uc = &OMP_CLAUSE_CHAIN (*uc);
}
else
pc = &OMP_CLAUSE_CHAIN (*pc);
*uc = NULL_TREE;
*pc = use_device_clauses;
stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_DATA,
OMP_CLAUSES (expr));
}
break;
case OMP_TEAMS:
stmt = gimple_build_omp_teams (body, OMP_CLAUSES (expr));

17
gcc/omp-low.c
View File

@ -1243,9 +1243,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR
&& !omp_is_reference (decl))
|| TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
install_var_field (decl, true, 3, ctx);
install_var_field (decl, true, 11, ctx);
else
install_var_field (decl, false, 3, ctx);
install_var_field (decl, false, 11, ctx);
if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
{
@ -11857,11 +11857,16 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
case OMP_CLAUSE_IS_DEVICE_PTR:
ovar = OMP_CLAUSE_DECL (c);
var = lookup_decl_in_outer_ctx (ovar, ctx);
x = build_sender_ref (ovar, ctx);
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR)
tkind = GOMP_MAP_USE_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;
{
tkind = GOMP_MAP_FIRSTPRIVATE_INT;
x = build_sender_ref (ovar, ctx);
}
type = TREE_TYPE (ovar);
if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR
&& !omp_is_reference (ovar))
@ -12032,7 +12037,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
case OMP_CLAUSE_IS_DEVICE_PTR:
var = OMP_CLAUSE_DECL (c);
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR)
x = build_sender_ref (var, ctx);
x = build_sender_ref ((splay_tree_key) &DECL_UID (var), ctx);
else
x = build_receiver_ref (var, false, ctx);
if (is_variable_sized (var))

9
libgomp/ChangeLog
View File

@ -1,3 +1,12 @@
2019-08-08 Jakub Jelinek <jakub@redhat.com>
* target.c (gomp_map_vars_internal): For GOMP_MAP_USE_DEVICE_PTR
perform the lookup in the first loop only if !not_found_cnt, otherwise
perform lookups for it in the second loop guarded with
if (not_found_cnt || has_firstprivate).
* testsuite/libgomp.c/target-37.c: New test.
* testsuite/libgomp.c++/target-22.C: New test.
2019-08-07 Jakub Jelinek <jakub@redhat.com>
* testsuite/libgomp.c/target-18.c (struct S): New type.

37
libgomp/target.c
View File

@ -580,20 +580,12 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
}
else if ((kind & typemask) == GOMP_MAP_USE_DEVICE_PTR)
{
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)
{
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].key = NULL;
tgt->list[i].offset = ~(uintptr_t) 0;
if (!not_found_cnt)
{
}
else
tgt->list[i].offset = 0;
continue;
}
else if ((kind & typemask) == GOMP_MAP_STRUCT)
@ -791,9 +783,26 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
tgt_size += len;
continue;
case GOMP_MAP_FIRSTPRIVATE_INT:
case GOMP_MAP_USE_DEVICE_PTR:
case GOMP_MAP_ZERO_LEN_ARRAY_SECTION:
continue;
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)
{
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;
}
continue;
case GOMP_MAP_STRUCT:
first = i + 1;
last = i + sizes[i];

99
libgomp/testsuite/libgomp.c++/target-22.C Normal file
View File

@ -0,0 +1,99 @@
extern "C" void abort (void);
struct S { int e, f; };
void
foo (int *&p, int (&s)[5], int &t, S &u, int n)
{
int a[4] = { 7, 8, 9, 10 }, b[n], c[3] = { 20, 21, 22 };
int *r = a + 1, *q = p - 1, i, err;
int v = 27;
S w = { 28, 29 };
for (i = 0; i < n; i++)
b[i] = 9 + i;
#pragma omp target data map(to:a) use_device_ptr(r) map(from:err)
#pragma omp target is_device_ptr(r) private(i) map(from:err)
{
err = 0;
for (i = 0; i < 4; i++)
if (r[i - 1] != 7 + i)
err = 1;
}
if (err)
abort ();
#pragma omp target data use_device_ptr(p) map(from:err) map(to:q[:4])
#pragma omp target is_device_ptr(p) private(i) map(from:err)
{
err = 0;
for (i = 0; i < 4; i++)
if (p[i - 1] != i)
err = 1;
}
if (err)
abort ();
#pragma omp target data map(to:b) use_device_addr(b) map(from:err)
#pragma omp target is_device_ptr(b) private(i) map(from:err)
{
err = 0;
for (i = 0; i < n; i++)
if (b[i] != 9 + i)
err = 1;
}
if (err)
abort ();
#pragma omp target data use_device_addr(c) map(to:c) map(from:err)
#pragma omp target is_device_ptr(c) private(i) map(from:err)
{
err = 0;
for (i = 0; i < 3; i++)
if (c[i] != 20 + i)
err = 1;
}
if (err)
abort ();
#pragma omp target data map(to:s[:5]) use_device_addr(s) map(from:err)
#pragma omp target is_device_ptr(s) private(i) map(from:err)
{
err = 0;
for (i = 0; i < 5; i++)
if (s[i] != 17 + i)
err = 1;
}
if (err)
abort ();
#pragma omp target data use_device_addr (v) map(to: v) map(to:u) use_device_addr (u) map(from:err)
{
int *z = &v;
S *x = &u;
#pragma omp target is_device_ptr (z, x) map(from:err)
{
err = 0;
if (*z != 27 || x->e != 25 || x->f != 26)
err = 1;
}
}
if (err)
abort ();
#pragma omp target data map(to: t) use_device_addr (t, w) map (to: w) map(from:err)
{
int *z = &t;
S *x = &w;
#pragma omp target is_device_ptr (z) is_device_ptr (x) map(from:err)
{
err = 0;
if (*z != 24 || x->e != 28 || x->f != 29)
err = 1;
}
}
if (err)
abort ();
}
int
main ()
{
int a[4] = { 0, 1, 2, 3 }, b[5] = { 17, 18, 19, 20, 21 };
int *p = a + 1;
int t = 24;
S u = { 25, 26 };
foo (p, b, t, u, 9);
}

71
libgomp/testsuite/libgomp.c/target-37.c Normal file
View File

@ -0,0 +1,71 @@
extern void abort (void);
struct S { int e, f; };
void
foo (int n)
{
int a[4] = { 0, 1, 2, 3 }, b[n], c = 4;
struct S d = { 5, 6 };
int *p = a + 1, i, err;
for (i = 0; i < n; i++)
b[i] = 9 + i;
#pragma omp target data use_device_ptr(p) map(from:err) map(to:a)
#pragma omp target is_device_ptr(p) private(i) map(from:err)
{
err = 0;
for (i = 0; i < 4; i++)
if (p[i - 1] != i)
err = 1;
}
if (err)
abort ();
for (i = 0; i < 4; i++)
a[i] = 23 + i;
#pragma omp target data map(to:a) use_device_addr(a) map(from:err)
#pragma omp target is_device_ptr(a) private(i) map(from:err)
{
err = 0;
for (i = 0; i < 4; i++)
if (a[i] != 23 + i)
err = 1;
}
if (err)
abort ();
#pragma omp target data use_device_addr(b) map(from:err) map(to:b)
#pragma omp target is_device_ptr(b) private(i) map(from:err)
{
err = 0;
for (i = 0; i < 4; i++)
if (b[i] != 9 + i)
err = 1;
}
if (err)
abort ();
#pragma omp target data map(to:c) use_device_addr(c) map(from:err)
{
int *q = &c;
#pragma omp target is_device_ptr(q) map(from:err)
{
err = *q != 4;
}
}
if (err)
abort ();
#pragma omp target data use_device_addr(d) map(to:d) map(from:err)
{
struct S *r = &d;
#pragma omp target is_device_ptr(r) map(from:err)
{
err = r->e != 5 || r->f != 6;
}
}
if (err)
abort ();
}
int
main ()
{
foo (9);
return 0;
}