On Thu, Sep 03, 2015 at 04:16:50PM +0200, Jakub Jelinek wrote:
> On Wed, Sep 02, 2015 at 05:58:54PM +0200, Jakub Jelinek wrote:
> > Here is the start of the async offloading support I've talked about,
> > but nowait is not supported on the library side yet, only depend clause
> > (and for that I haven't added a testcase yet).
>
> Added testcase revealed two (small) issues, here is a fix for that together
> with the testcase.
There has been a bug in the testcase (missing map(from:err) in 3 places),
which hid a problem that on target constructs with depend clause (what about
just nowait?) we have to avoid using GOMP_FIRSTPRIVATE_INT or copy value
into temporary and take temporary's address for GOMP_FIRSTPRIVATE unless
we can prove other tasks can't modify the value while waiting for
dependencies (if it is addressable or shared with other threads/tasks,
then we have to use GOMP_FIRSTPRIVATE with address of the real variable,
so that if other tasks change it, we pick up the right values).
2015-09-04 Jakub Jelinek <ja...@redhat.com>
* omp-low.c (lower_omp_target): If target has depend
clauses, avoid using GOMP_MAP_FIRSTPRIVATE_INT unless the
var is non-addressable and private in the current task.
Even for GOMP_MAP_FIRSTPRIVATE, if the var is non-addressable,
but shared or threadprivate, take address of the shared var
rather than initializing a temporary with the current value.
* testsuite/libgomp.c/target-25.c (main): Add missing
map(from: err) clauses to target constructs.
--- gcc/omp-low.c.jj 2015-09-03 16:36:31.000000000 +0200
+++ gcc/omp-low.c 2015-09-04 11:34:45.512416693 +0200
@@ -13236,6 +13236,7 @@ lower_omp_target (gimple_stmt_iterator *
location_t loc = gimple_location (stmt);
bool offloaded, data_region;
unsigned int map_cnt = 0;
+ bool has_depend = false;
offloaded = is_gimple_omp_offloaded (stmt);
switch (gimple_omp_target_kind (stmt))
@@ -13268,6 +13269,7 @@ lower_omp_target (gimple_stmt_iterator *
dep_bind = gimple_build_bind (NULL, NULL, make_node (BLOCK));
lower_depend_clauses (gimple_omp_target_clauses_ptr (stmt),
&dep_ilist, &dep_olist);
+ has_depend = true;
}
tgt_bind = NULL;
@@ -13719,9 +13721,44 @@ lower_omp_target (gimple_stmt_iterator *
type = TREE_TYPE (ovar);
if (is_reference (ovar))
type = TREE_TYPE (type);
+ bool use_firstprivate_int, force_addr;
+ use_firstprivate_int = false;
+ force_addr = false;
if ((INTEGRAL_TYPE_P (type)
- && TYPE_PRECISION (type) <= POINTER_SIZE)
+ && TYPE_PRECISION (type) <= POINTER_SIZE)
|| TREE_CODE (type) == POINTER_TYPE)
+ use_firstprivate_int = true;
+ if (has_depend)
+ {
+ if (is_reference (var))
+ use_firstprivate_int = false;
+ else if (is_gimple_reg (var))
+ {
+ if (DECL_HAS_VALUE_EXPR_P (var))
+ {
+ tree v = get_base_address (var);
+ if (DECL_P (v) && TREE_ADDRESSABLE (v))
+ {
+ use_firstprivate_int = false;
+ force_addr = true;
+ }
+ else
+ switch (TREE_CODE (v))
+ {
+ case INDIRECT_REF:
+ case MEM_REF:
+ use_firstprivate_int = false;
+ force_addr = true;
+ break;
+ default:
+ break;
+ }
+ }
+ }
+ else
+ use_firstprivate_int = false;
+ }
+ if (use_firstprivate_int)
{
tkind = GOMP_MAP_FIRSTPRIVATE_INT;
tree t = var;
@@ -13734,7 +13771,7 @@ lower_omp_target (gimple_stmt_iterator *
}
else if (is_reference (var))
gimplify_assign (x, var, &ilist);
- else if (is_gimple_reg (var))
+ else if (!force_addr && is_gimple_reg (var))
{
tree avar = create_tmp_var (TREE_TYPE (var));
mark_addressable (avar);
@@ -13867,9 +13904,40 @@ lower_omp_target (gimple_stmt_iterator *
type = TREE_TYPE (var);
if (is_reference (var))
type = TREE_TYPE (type);
+ bool use_firstprivate_int;
+ use_firstprivate_int = false;
if ((INTEGRAL_TYPE_P (type)
&& TYPE_PRECISION (type) <= POINTER_SIZE)
|| TREE_CODE (type) == POINTER_TYPE)
+ use_firstprivate_int = true;
+ if (has_depend)
+ {
+ tree v = lookup_decl_in_outer_ctx (var, ctx);
+ if (is_reference (v))
+ use_firstprivate_int = false;
+ else if (is_gimple_reg (v))
+ {
+ if (DECL_HAS_VALUE_EXPR_P (v))
+ {
+ v = get_base_address (v);
+ if (DECL_P (v) && TREE_ADDRESSABLE (v))
+ use_firstprivate_int = false;
+ else
+ switch (TREE_CODE (v))
+ {
+ case INDIRECT_REF:
+ case MEM_REF:
+ use_firstprivate_int = false;
+ break;
+ default:
+ break;
+ }
+ }
+ }
+ else
+ use_firstprivate_int = false;
+ }
+ if (use_firstprivate_int)
{
x = build_receiver_ref (var, false, ctx);
if (TREE_CODE (type) != POINTER_TYPE)
--- libgomp/testsuite/libgomp.c/target-25.c.jj 2015-09-04 10:41:52.371881507
+0200
+++ libgomp/testsuite/libgomp.c/target-25.c 2015-09-04 10:39:16.000000000
+0200
@@ -23,7 +23,7 @@ main ()
usleep (7000);
z = 3;
}
- #pragma omp target map(tofrom: x) firstprivate (y) depend(inout: x, z)
+ #pragma omp target map(tofrom: x) map(from: err) firstprivate (y)
depend(inout: x, z)
err = (x != 1 || y != 2 || z != 3);
if (err)
abort ();
@@ -44,7 +44,7 @@ main ()
}
#pragma omp target enter data nowait map (to: w)
#pragma omp target enter data depend (inout: x, z) map (to: x, y, z)
- #pragma omp target map (alloc: x, y, z)
+ #pragma omp target map (alloc: x, y, z) map(from: err)
{
err = (x != 4 || y != 5 || z != 6);
x = 7;
@@ -54,7 +54,7 @@ main ()
if (err)
abort ();
#pragma omp taskwait
- #pragma omp target map (alloc: w)
+ #pragma omp target map (alloc: w) map(from: err)
{
err = w != 7;
w = 17;
Jakub
