* OpenACC Firstprivate
@ 2015-11-07 13:50 Nathan Sidwell
2015-11-09 13:46 ` Jakub Jelinek
0 siblings, 1 reply; 11+ messages in thread
From: Nathan Sidwell @ 2015-11-07 13:50 UTC (permalink / raw)
To: Jakub Jelinek; +Cc: GCC Patches, Cesar Philippidis
[-- Attachment #1: Type: text/plain, Size: 1520 bytes --]
Jakub,
this patch implements firstprivate support for openacc. This is pretty straight
forwards -- they're just regular auto variables, but with an initialization
value from the host.
The gimplify.c implementation is somewhat different to gomp4 branch, as I've
added new bits to enum omp_region_type, rather than add 2 new fields to
omp_region_ctx. The new enums use bits already defined in omp_region_type:
+ ORT_ACC = 0x40, /* An OpenACC region. */
+ ORT_ACC_DATA = ORT_ACC | ORT_TARGET_DATA, /* Data construct. */
+ ORT_ACC_PARALLEL = ORT_ACC | ORT_TARGET, /* Parallel construct */
+ ORT_ACC_KERNELS = ORT_ACC | ORT_TARGET | 0x80, /* Kernels construct. */
On gomp4 we were already setting those bits, but then setting the new fields to
indicate 'openacc'. Many places in gimplify.c where we check for '==
ORT_TARGET_DATA' or ORT_TARGET get changed to '& ORT_TARGET_DATA' etc.
On gomp4 for things like an openacc loop we were setting ORT_WORKSHARE, so
nearly all checks for == ORT_WORKSHARE get an additional '|| X == ORT_ACC'.
Although this patch doesn't make use of the difference between ORT_ACC_KERNELS
and ORT_ACC_PARALLEL, the default handling patch will -- they have different
behaviours.
I think the gimpify.c changes are then obvious from that, but let me know.
in omp-low the changes are to remove 'sorry' and build the initializer exprs in
lower_omp_target.
As you can see this fixes a few xfails.
I'll post the default handling patch, which is much more localized.
nathan
[-- Attachment #2: trunk-firstprivate-1106.patch --]
[-- Type: text/x-patch, Size: 21221 bytes --]
2015-11-06 Nathan Sidwell <nathan@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
gcc/
* gcc/gimplify.c (enum omp_region_type): Add ORT_ACC,
ORT_ACC_DATA, ORT_ACC_PARALLEL, ORT_ACC_KERNELS. Adjust ORT_NONE.
(new_omp_context): Initialize all fields.
(gimple_add_tmp_var): Add ORT_ACC checks.
(gimplify_var_or_parm_decl): Likewise.
(omp_firstprivatize_variable): Likewise. Use ORT_TARGET_DATA as a
mask.
(omp_add_variable): Look in outer contexts for openacc and allow
reductions with other sharing. Add ORT_ACC and ORT_TARGET_DATA
checks.
(omp_notice_variable, omp_is_private, omp_check_private): Add
ORT_ACC checks.
(gimplify_scan_omp_clauses: Treat ORT_ACC as ORT_WORKSHARE.
Permit private openacc reductions.
(gimplify_oacc_cache): Specify ORT_ACC.
(gimplify_omp_workshare): Adjust OpenACC region types.
(gimplify_omp_target_update): Likewise.
* gcc/omp-low.c (scan_sharing_clauses): Remove Openacc
firstprivate sorry.
(lower-rec_input_clauses): Don't handle openacc firstprivate
references here.
(lower_omp_target): Emit initializers for openacc firstprivate vars.
gcc/testsuite/
* gfortran.dg/goacc/private-3.f95: Remove xfail.
* gfortran.dg/goacc/combined_loop.f90: Remove xfail.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c: Remove xfail.
* testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c: Remove xfail.
* testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c: New.
Index: gcc/gimplify.c
===================================================================
--- gcc/gimplify.c (revision 229892)
+++ gcc/gimplify.c (working copy)
@@ -108,9 +108,15 @@ enum omp_region_type
/* Data region with offloading. */
ORT_TARGET = 32,
ORT_COMBINED_TARGET = 33,
+
+ ORT_ACC = 0x40, /* An OpenACC region. */
+ ORT_ACC_DATA = ORT_ACC | ORT_TARGET_DATA, /* Data construct. */
+ ORT_ACC_PARALLEL = ORT_ACC | ORT_TARGET, /* Parallel construct */
+ ORT_ACC_KERNELS = ORT_ACC | ORT_TARGET | 0x80, /* Kernels construct. */
+
/* Dummy OpenMP region, used to disable expansion of
DECL_VALUE_EXPRs in taskloop pre body. */
- ORT_NONE = 64
+ ORT_NONE = 0x100
};
/* Gimplify hashtable helper. */
@@ -377,6 +383,12 @@ new_omp_context (enum omp_region_type re
else
c->default_kind = OMP_CLAUSE_DEFAULT_UNSPECIFIED;
+ c->combined_loop = false;
+ c->distribute = false;
+ c->target_map_scalars_firstprivate = false;
+ c->target_map_pointers_as_0len_arrays = false;
+ c->target_firstprivatize_array_bases = false;
+
return c;
}
@@ -689,7 +701,8 @@ gimple_add_tmp_var (tree tmp)
struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
while (ctx
&& (ctx->region_type == ORT_WORKSHARE
- || ctx->region_type == ORT_SIMD))
+ || ctx->region_type == ORT_SIMD
+ || ctx->region_type == ORT_ACC))
ctx = ctx->outer_context;
if (ctx)
omp_add_variable (ctx, tmp, GOVD_LOCAL | GOVD_SEEN);
@@ -1804,7 +1817,8 @@ gimplify_var_or_parm_decl (tree *expr_p)
struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
while (ctx
&& (ctx->region_type == ORT_WORKSHARE
- || ctx->region_type == ORT_SIMD))
+ || ctx->region_type == ORT_SIMD
+ || ctx->region_type == ORT_ACC))
ctx = ctx->outer_context;
if (!ctx && !nonlocal_vlas->add (decl))
{
@@ -5579,7 +5593,8 @@ omp_firstprivatize_variable (struct gimp
}
else if (ctx->region_type != ORT_WORKSHARE
&& ctx->region_type != ORT_SIMD
- && ctx->region_type != ORT_TARGET_DATA)
+ && ctx->region_type != ORT_ACC
+ && !(ctx->region_type & ORT_TARGET_DATA))
omp_add_variable (ctx, decl, GOVD_FIRSTPRIVATE);
ctx = ctx->outer_context;
@@ -5667,11 +5682,13 @@ omp_add_variable (struct gimplify_omp_ct
/* We shouldn't be re-adding the decl with the same data
sharing class. */
gcc_assert ((n->value & GOVD_DATA_SHARE_CLASS & flags) == 0);
- /* The only combination of data sharing classes we should see is
- FIRSTPRIVATE and LASTPRIVATE. */
nflags = n->value | flags;
- gcc_assert ((nflags & GOVD_DATA_SHARE_CLASS)
- == (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE)
+ /* The only combination of data sharing classes we should see is
+ FIRSTPRIVATE and LASTPRIVATE. However, OpenACC permits
+ reduction variables to be used in data sharing clauses. */
+ gcc_assert ((ctx->region_type & ORT_ACC) != 0
+ || ((nflags & GOVD_DATA_SHARE_CLASS)
+ == (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE))
|| (flags & GOVD_DATA_SHARE_CLASS) == 0);
n->value = nflags;
return;
@@ -5968,20 +5985,47 @@ omp_notice_variable (struct gimplify_omp
else if (is_scalar)
nflags |= GOVD_FIRSTPRIVATE;
}
- tree type = TREE_TYPE (decl);
- if (nflags == flags
- && gimplify_omp_ctxp->target_firstprivatize_array_bases
- && lang_hooks.decls.omp_privatize_by_reference (decl))
- type = TREE_TYPE (type);
- if (nflags == flags
- && !lang_hooks.types.omp_mappable_type (type))
- {
- error ("%qD referenced in target region does not have "
- "a mappable type", decl);
- nflags |= GOVD_MAP | GOVD_EXPLICIT;
+
+ /* OpenMP doesn't look in outer contexts to find an
+ enclosing data clause. */
+ struct gimplify_omp_ctx *octx = ctx->outer_context;
+ if ((ctx->region_type & ORT_ACC) && octx)
+ {
+ omp_notice_variable (octx, decl, in_code);
+
+ for (; octx; octx = octx->outer_context)
+ {
+ if (!(octx->region_type & (ORT_TARGET_DATA | ORT_TARGET)))
+ break;
+ splay_tree_node n2
+ = splay_tree_lookup (octx->variables,
+ (splay_tree_key) decl);
+ if (n2)
+ {
+ nflags |= GOVD_MAP;
+ goto found_outer;
+ }
+ }
}
- else if (nflags == flags)
- nflags |= GOVD_MAP;
+
+ {
+ tree type = TREE_TYPE (decl);
+
+ if (nflags == flags
+ && gimplify_omp_ctxp->target_firstprivatize_array_bases
+ && lang_hooks.decls.omp_privatize_by_reference (decl))
+ type = TREE_TYPE (type);
+ if (nflags == flags
+ && !lang_hooks.types.omp_mappable_type (type))
+ {
+ error ("%qD referenced in target region does not have "
+ "a mappable type", decl);
+ nflags |= GOVD_MAP | GOVD_EXPLICIT;
+ }
+ else if (nflags == flags)
+ nflags |= GOVD_MAP;
+ }
+ found_outer:
omp_add_variable (ctx, decl, nflags);
}
else
@@ -5998,7 +6042,8 @@ omp_notice_variable (struct gimplify_omp
{
if (ctx->region_type == ORT_WORKSHARE
|| ctx->region_type == ORT_SIMD
- || ctx->region_type == ORT_TARGET_DATA)
+ || ctx->region_type == ORT_ACC
+ || (ctx->region_type & ORT_TARGET_DATA) != 0)
goto do_outer;
flags = omp_default_clause (ctx, decl, in_code, flags);
@@ -6112,7 +6157,8 @@ omp_is_private (struct gimplify_omp_ctx
}
if (ctx->region_type != ORT_WORKSHARE
- && ctx->region_type != ORT_SIMD)
+ && ctx->region_type != ORT_SIMD
+ && ctx->region_type != ORT_ACC)
return false;
else if (ctx->outer_context)
return omp_is_private (ctx->outer_context, decl, simd);
@@ -6168,7 +6214,8 @@ omp_check_private (struct gimplify_omp_c
}
}
while (ctx->region_type == ORT_WORKSHARE
- || ctx->region_type == ORT_SIMD);
+ || ctx->region_type == ORT_SIMD
+ || ctx->region_type == ORT_ACC);
return false;
}
@@ -6311,7 +6358,8 @@ gimplify_scan_omp_clauses (tree *list_p,
omp_notice_variable (outer_ctx->outer_context, decl, true);
}
else if (outer_ctx
- && outer_ctx->region_type == ORT_WORKSHARE
+ && (outer_ctx->region_type == ORT_WORKSHARE
+ || outer_ctx->region_type == ORT_ACC)
&& outer_ctx->combined_loop
&& splay_tree_lookup (outer_ctx->variables,
(splay_tree_key) decl) == NULL
@@ -6335,7 +6383,9 @@ gimplify_scan_omp_clauses (tree *list_p,
goto do_add;
case OMP_CLAUSE_REDUCTION:
flags = GOVD_REDUCTION | GOVD_SEEN | GOVD_EXPLICIT;
- check_non_private = "reduction";
+ /* OpenACC permits reductions on private variables. */
+ if (!(region_type & ORT_ACC))
+ check_non_private = "reduction";
decl = OMP_CLAUSE_DECL (c);
if (TREE_CODE (decl) == MEM_REF)
{
@@ -7703,7 +7753,7 @@ gimplify_oacc_cache (tree *expr_p, gimpl
{
tree expr = *expr_p;
- gimplify_scan_omp_clauses (&OACC_CACHE_CLAUSES (expr), pre_p, ORT_WORKSHARE,
+ gimplify_scan_omp_clauses (&OACC_CACHE_CLAUSES (expr), pre_p, ORT_ACC,
OACC_CACHE);
gimplify_adjust_omp_clauses (pre_p, &OACC_CACHE_CLAUSES (expr), OACC_CACHE);
@@ -7832,7 +7882,9 @@ gimplify_omp_for (tree *expr_p, gimple_s
case OMP_FOR:
case CILK_FOR:
case OMP_DISTRIBUTE:
+ break;
case OACC_LOOP:
+ ort = ORT_ACC;
break;
case OMP_TASKLOOP:
if (find_omp_clause (OMP_FOR_CLAUSES (for_stmt), OMP_CLAUSE_UNTIED))
@@ -8894,10 +8946,14 @@ gimplify_omp_workshare (tree *expr_p, gi
ort = OMP_TARGET_COMBINED (expr) ? ORT_COMBINED_TARGET : ORT_TARGET;
break;
case OACC_KERNELS:
+ ort = ORT_ACC_KERNELS;
+ break;
case OACC_PARALLEL:
- ort = ORT_TARGET;
+ ort = ORT_ACC_PARALLEL;
break;
case OACC_DATA:
+ ort = ORT_ACC_DATA;
+ break;
case OMP_TARGET_DATA:
ort = ORT_TARGET_DATA;
break;
@@ -8919,7 +8975,7 @@ gimplify_omp_workshare (tree *expr_p, gi
pop_gimplify_context (g);
else
pop_gimplify_context (NULL);
- if (ort == ORT_TARGET_DATA)
+ if ((ort & ORT_TARGET_DATA) != 0)
{
enum built_in_function end_ix;
switch (TREE_CODE (expr))
@@ -8994,17 +9050,18 @@ gimplify_omp_target_update (tree *expr_p
tree expr = *expr_p;
int kind;
gomp_target *stmt;
+ enum omp_region_type ort = ORT_WORKSHARE;
switch (TREE_CODE (expr))
{
case OACC_ENTER_DATA:
- kind = GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA;
- break;
case OACC_EXIT_DATA:
kind = GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA;
+ ort = ORT_ACC;
break;
case OACC_UPDATE:
kind = GF_OMP_TARGET_KIND_OACC_UPDATE;
+ ort = ORT_ACC;
break;
case OMP_TARGET_UPDATE:
kind = GF_OMP_TARGET_KIND_UPDATE;
@@ -9019,7 +9076,7 @@ gimplify_omp_target_update (tree *expr_p
gcc_unreachable ();
}
gimplify_scan_omp_clauses (&OMP_STANDALONE_CLAUSES (expr), pre_p,
- ORT_WORKSHARE, TREE_CODE (expr));
+ ort, TREE_CODE (expr));
gimplify_adjust_omp_clauses (pre_p, &OMP_STANDALONE_CLAUSES (expr),
TREE_CODE (expr));
stmt = gimple_build_omp_target (NULL, kind, OMP_STANDALONE_CLAUSES (expr));
Index: gcc/omp-low.c
===================================================================
--- gcc/omp-low.c (revision 229892)
+++ gcc/omp-low.c (working copy)
@@ -1896,12 +1896,6 @@ scan_sharing_clauses (tree clauses, omp_
/* FALLTHRU */
case OMP_CLAUSE_FIRSTPRIVATE:
- if (is_gimple_omp_oacc (ctx->stmt))
- {
- sorry ("clause not supported yet");
- break;
- }
- /* FALLTHRU */
case OMP_CLAUSE_LINEAR:
decl = OMP_CLAUSE_DECL (c);
do_private:
@@ -2167,12 +2161,6 @@ scan_sharing_clauses (tree clauses, omp_
/* FALLTHRU */
case OMP_CLAUSE_FIRSTPRIVATE:
- if (is_gimple_omp_oacc (ctx->stmt))
- {
- sorry ("clause not supported yet");
- break;
- }
- /* FALLTHRU */
case OMP_CLAUSE_PRIVATE:
case OMP_CLAUSE_LINEAR:
case OMP_CLAUSE_IS_DEVICE_PTR:
@@ -4684,7 +4672,7 @@ lower_rec_input_clauses (tree clauses, g
gimplify_assign (ptr, x, ilist);
}
}
- else if (is_reference (var))
+ else if (is_reference (var) && !is_oacc_parallel (ctx))
{
/* For references that are being privatized for Fortran,
allocate new backing storage for the new pointer
@@ -14878,7 +14866,7 @@ lower_omp_target (gimple_stmt_iterator *
tree child_fn, t, c;
gomp_target *stmt = as_a <gomp_target *> (gsi_stmt (*gsi_p));
gbind *tgt_bind, *bind, *dep_bind = NULL;
- gimple_seq tgt_body, olist, ilist, new_body;
+ gimple_seq tgt_body, olist, ilist, fplist, new_body;
location_t loc = gimple_location (stmt);
bool offloaded, data_region;
unsigned int map_cnt = 0;
@@ -14930,6 +14918,7 @@ lower_omp_target (gimple_stmt_iterator *
child_fn = ctx->cb.dst_fn;
push_gimplify_context ();
+ fplist = NULL;
for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
switch (OMP_CLAUSE_CODE (c))
@@ -14974,6 +14963,7 @@ lower_omp_target (gimple_stmt_iterator *
/* FALLTHRU */
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
+ oacc_firstprivate:
var = OMP_CLAUSE_DECL (c);
if (!DECL_P (var))
{
@@ -14996,6 +14986,7 @@ lower_omp_target (gimple_stmt_iterator *
}
if (offloaded
+ && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
{
@@ -15024,17 +15015,40 @@ lower_omp_target (gimple_stmt_iterator *
x = build_receiver_ref (var, true, ctx);
tree new_var = lookup_decl (var, ctx);
- if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
&& TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
x = build_simple_mem_ref (x);
- SET_DECL_VALUE_EXPR (new_var, x);
- DECL_HAS_VALUE_EXPR_P (new_var) = 1;
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
+ {
+ gcc_assert (is_gimple_omp_oacc (ctx->stmt));
+ if (is_reference (new_var))
+ {
+ /* Create a local object to hold the instance
+ value. */
+ tree inst = create_tmp_var
+ (TREE_TYPE (TREE_TYPE (new_var)),
+ IDENTIFIER_POINTER (DECL_NAME (new_var)));
+ gimplify_assign (inst, fold_indirect_ref (x), &fplist);
+ x = build_fold_addr_expr (inst);
+ }
+ gimplify_assign (new_var, x, &fplist);
+ }
+ else if (DECL_P (new_var))
+ {
+ SET_DECL_VALUE_EXPR (new_var, x);
+ DECL_HAS_VALUE_EXPR_P (new_var) = 1;
+ }
+ else
+ gcc_unreachable ();
}
map_cnt++;
break;
case OMP_CLAUSE_FIRSTPRIVATE:
+ if (is_oacc_parallel (ctx))
+ goto oacc_firstprivate;
map_cnt++;
var = OMP_CLAUSE_DECL (c);
if (!is_reference (var)
@@ -15059,6 +15073,8 @@ lower_omp_target (gimple_stmt_iterator *
break;
case OMP_CLAUSE_PRIVATE:
+ if (is_gimple_omp_oacc (ctx->stmt))
+ break;
var = OMP_CLAUSE_DECL (c);
if (is_variable_sized (var))
{
@@ -15162,9 +15178,11 @@ lower_omp_target (gimple_stmt_iterator *
default:
break;
+
case OMP_CLAUSE_MAP:
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
+ oacc_firstprivate_map:
nc = c;
ovar = OMP_CLAUSE_DECL (c);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
@@ -15215,9 +15233,9 @@ lower_omp_target (gimple_stmt_iterator *
x = build_sender_ref (ovar, ctx);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
- && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
- && TREE_CODE (TREE_TYPE (ovar)) == ARRAY_TYPE)
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
+ && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
+ && TREE_CODE (TREE_TYPE (ovar)) == ARRAY_TYPE)
{
gcc_assert (offloaded);
tree avar
@@ -15228,6 +15246,15 @@ lower_omp_target (gimple_stmt_iterator *
avar = build_fold_addr_expr (avar);
gimplify_assign (x, avar, &ilist);
}
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
+ {
+ gcc_assert (is_gimple_omp_oacc (ctx->stmt));
+ if (!is_reference (var))
+ var = build_fold_addr_expr (var);
+ else
+ talign = TYPE_ALIGN_UNIT (TREE_TYPE (TREE_TYPE (ovar)));
+ gimplify_assign (x, var, &ilist);
+ }
else if (is_gimple_reg (var))
{
gcc_assert (offloaded);
@@ -15256,7 +15283,17 @@ lower_omp_target (gimple_stmt_iterator *
gimplify_assign (x, var, &ilist);
}
}
- s = OMP_CLAUSE_SIZE (c);
+ s = NULL_TREE;
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
+ {
+ gcc_checking_assert (is_gimple_omp_oacc (ctx->stmt));
+ s = TREE_TYPE (ovar);
+ if (TREE_CODE (s) == REFERENCE_TYPE)
+ s = TREE_TYPE (s);
+ s = TYPE_SIZE_UNIT (s);
+ }
+ else
+ s = OMP_CLAUSE_SIZE (c);
if (s == NULL_TREE)
s = TYPE_SIZE_UNIT (TREE_TYPE (ovar));
s = fold_convert (size_type_node, s);
@@ -15297,6 +15334,11 @@ lower_omp_target (gimple_stmt_iterator *
tkind_zero = tkind;
}
break;
+ case OMP_CLAUSE_FIRSTPRIVATE:
+ gcc_checking_assert (is_gimple_omp_oacc (ctx->stmt));
+ tkind = GOMP_MAP_TO;
+ tkind_zero = tkind;
+ break;
case OMP_CLAUSE_TO:
tkind = GOMP_MAP_TO;
tkind_zero = tkind;
@@ -15336,6 +15378,8 @@ lower_omp_target (gimple_stmt_iterator *
break;
case OMP_CLAUSE_FIRSTPRIVATE:
+ if (is_oacc_parallel (ctx))
+ goto oacc_firstprivate_map;
ovar = OMP_CLAUSE_DECL (c);
if (is_reference (ovar))
talign = TYPE_ALIGN_UNIT (TREE_TYPE (TREE_TYPE (ovar)));
@@ -15510,6 +15554,7 @@ lower_omp_target (gimple_stmt_iterator *
gimple_seq_add_stmt (&new_body,
gimple_build_assign (ctx->receiver_decl, t));
}
+ gimple_seq_add_seq (&new_body, fplist);
if (offloaded || data_region)
{
@@ -15521,6 +15566,8 @@ lower_omp_target (gimple_stmt_iterator *
default:
break;
case OMP_CLAUSE_FIRSTPRIVATE:
+ if (is_gimple_omp_oacc (ctx->stmt))
+ break;
var = OMP_CLAUSE_DECL (c);
if (is_reference (var)
|| is_gimple_reg_type (TREE_TYPE (var)))
@@ -15606,6 +15653,8 @@ lower_omp_target (gimple_stmt_iterator *
}
break;
case OMP_CLAUSE_PRIVATE:
+ if (is_gimple_omp_oacc (ctx->stmt))
+ break;
var = OMP_CLAUSE_DECL (c);
if (is_reference (var))
{
@@ -15694,7 +15743,7 @@ lower_omp_target (gimple_stmt_iterator *
/* Handle GOMP_MAP_FIRSTPRIVATE_{POINTER,REFERENCE} in second pass,
so that firstprivate vars holding OMP_CLAUSE_SIZE if needed
are already handled. */
- for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
+ for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
switch (OMP_CLAUSE_CODE (c))
{
tree var;
Index: gcc/testsuite/gfortran.dg/goacc/private-3.f95
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/private-3.f95 (revision 229864)
+++ gcc/testsuite/gfortran.dg/goacc/private-3.f95 (working copy)
@@ -1,6 +1,4 @@
! { dg-do compile }
-! <http://news.gmane.org/find-root.php?message_id=%3C563B78B5.5090506%40acm.org%3E>
-! { dg-xfail-if "TODO" { *-*-* } }
! test for private variables in a reduction clause
Index: gcc/testsuite/gfortran.dg/goacc/combined_loop.f90
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/combined_loop.f90 (revision 229864)
+++ gcc/testsuite/gfortran.dg/goacc/combined_loop.f90 (working copy)
@@ -1,6 +1,4 @@
! { dg-do compile }
-! <http://news.gmane.org/find-root.php?message_id=%3C563B78B5.5090506%40acm.org%3E>
-! { dg-xfail-if "TODO" { *-*-* } }
!
! PR fortran/64726
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c (revision 229852)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c (working copy)
@@ -1,7 +1,5 @@
/* { dg-do run } */
/* { dg-additional-options "-O2" */
-/* <http://news.gmane.org/find-root.php?message_id=%3C563B78B5.5090506%40acm.org%3E>
- { dg-xfail-if "TODO" { *-*-* } } */
#include <stdio.h>
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c (revision 229852)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c (working copy)
@@ -1,7 +1,5 @@
/* { dg-do run } */
/* { dg-additional-options "-O2" */
-/* <http://news.gmane.org/find-root.php?message_id=%3C563B78B5.5090506%40acm.org%3E>
- { dg-xfail-if "TODO" { *-*-* } } */
#include <stdio.h>
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c (revision 0)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c (working copy)
@@ -0,0 +1,41 @@
+/* { dg-do run } */
+
+#include <openacc.h>
+
+int main ()
+{
+ int ok = 1;
+ int val = 2;
+ int ary[32];
+ int ondev = 0;
+
+ for (int i = 0; i < 32; i++)
+ ary[i] = ~0;
+
+#pragma acc parallel num_gangs (32) copy (ok) firstprivate (val) copy(ary, ondev)
+ {
+ ondev = acc_on_device (acc_device_not_host);
+#pragma acc loop gang(static:1)
+ for (unsigned i = 0; i < 32; i++)
+ {
+ if (val != 2)
+ ok = 0;
+ val += i;
+ ary[i] = val;
+ }
+ }
+
+ if (ondev)
+ {
+ if (!ok)
+ return 1;
+ if (val != 2)
+ return 1;
+
+ for (int i = 0; i < 32; i++)
+ if (ary[i] != 2 + i)
+ return 1;
+ }
+
+ return 0;
+}
^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: OpenACC Firstprivate
2015-11-07 13:50 OpenACC Firstprivate Nathan Sidwell
@ 2015-11-09 13:46 ` Jakub Jelinek
2015-11-09 13:59 ` Nathan Sidwell
0 siblings, 1 reply; 11+ messages in thread
From: Jakub Jelinek @ 2015-11-09 13:46 UTC (permalink / raw)
To: Nathan Sidwell; +Cc: GCC Patches, Cesar Philippidis
On Sat, Nov 07, 2015 at 08:50:28AM -0500, Nathan Sidwell wrote:
> Index: gcc/gimplify.c
> ===================================================================
> --- gcc/gimplify.c (revision 229892)
> +++ gcc/gimplify.c (working copy)
> @@ -108,9 +108,15 @@ enum omp_region_type
> /* Data region with offloading. */
> ORT_TARGET = 32,
> ORT_COMBINED_TARGET = 33,
> +
> + ORT_ACC = 0x40, /* An OpenACC region. */
> + ORT_ACC_DATA = ORT_ACC | ORT_TARGET_DATA, /* Data construct. */
> + ORT_ACC_PARALLEL = ORT_ACC | ORT_TARGET, /* Parallel construct */
> + ORT_ACC_KERNELS = ORT_ACC | ORT_TARGET | 0x80, /* Kernels construct. */
> +
> /* Dummy OpenMP region, used to disable expansion of
> DECL_VALUE_EXPRs in taskloop pre body. */
> - ORT_NONE = 64
> + ORT_NONE = 0x100
> };
If you want to switch to hexadecimal, you should change all values
in the enum to hexadecimal for consistency.
>
> /* Gimplify hashtable helper. */
> @@ -377,6 +383,12 @@ new_omp_context (enum omp_region_type re
> else
> c->default_kind = OMP_CLAUSE_DEFAULT_UNSPECIFIED;
>
> + c->combined_loop = false;
> + c->distribute = false;
> + c->target_map_scalars_firstprivate = false;
> + c->target_map_pointers_as_0len_arrays = false;
> + c->target_firstprivatize_array_bases = false;
Why this? c is XCNEW allocated, so zero initialized.
> @@ -5667,11 +5682,13 @@ omp_add_variable (struct gimplify_omp_ct
> /* We shouldn't be re-adding the decl with the same data
> sharing class. */
> gcc_assert ((n->value & GOVD_DATA_SHARE_CLASS & flags) == 0);
> - /* The only combination of data sharing classes we should see is
> - FIRSTPRIVATE and LASTPRIVATE. */
> nflags = n->value | flags;
> - gcc_assert ((nflags & GOVD_DATA_SHARE_CLASS)
> - == (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE)
> + /* The only combination of data sharing classes we should see is
> + FIRSTPRIVATE and LASTPRIVATE. However, OpenACC permits
> + reduction variables to be used in data sharing clauses. */
> + gcc_assert ((ctx->region_type & ORT_ACC) != 0
> + || ((nflags & GOVD_DATA_SHARE_CLASS)
> + == (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE))
> || (flags & GOVD_DATA_SHARE_CLASS) == 0);
Are you sure you want to give up on any kind of consistency checks for
OpenACC? If only reduction is special on OpenACC, perhaps you could tweak
the assert for that instead? Something that can be done incrementally of
course.
> +
> + /* OpenMP doesn't look in outer contexts to find an
> + enclosing data clause. */
I'm puzzled by the comment. OpenMP does look in outer context for clauses
that need that (pretty much all closes but private), that is the do_outer:
recursion in omp_notice_variable. Say for firstprivate in order to copy (or
copy construct) the private variable one needs the access to the outer
context's var etc.).
So perhaps it would help to document what you are doing here for OpenACC and
why.
> + struct gimplify_omp_ctx *octx = ctx->outer_context;
> + if ((ctx->region_type & ORT_ACC) && octx)
> + {
> + omp_notice_variable (octx, decl, in_code);
> +
> + for (; octx; octx = octx->outer_context)
> + {
> + if (!(octx->region_type & (ORT_TARGET_DATA | ORT_TARGET)))
> + break;
> + splay_tree_node n2
> + = splay_tree_lookup (octx->variables,
> + (splay_tree_key) decl);
> + if (n2)
> + {
> + nflags |= GOVD_MAP;
> + goto found_outer;
> + }
> + }
> }
> - else if (nflags == flags)
> - nflags |= GOVD_MAP;
> +
The main issue I have is with the omp-low.c changes.
I see:
"2.5.9
private clause
The private clause is allowed on the parallel construct; it declares that a copy of each
item on the list will be created for each parallel gang.
2.5.10
firstprivate clause
The firstprivate clause is allowed on the parallel construct; it declares that a copy
of each item on the list will be created for each parallel gang, and that the copy will be
initialized with the value of that item on the host when the parallel construct is
encountered."
but looking at what you actually emit looks like standard present_copyin
clause I think with a private variable defined in the region where the
value of the present_copyin mapped variable is assigned to the private one.
This I'm afraid performs often two copies rather than just one (one to copy
the host value to the present_copyin mapped value, another one in the
region), but more importantly, if the var is already mapped, you could
initialize the private var with old data.
Say
int arr[64];
// initialize arr
#pragma acc data copyin (arr)
{
// modify arr on the host
# pragma acc parallel firstprivate (arr)
{
...
}
}
Is that really what you want? If not, any reason not to implement
GOMP_MAP_FIRSTPRIVATE and GOMP_MAP_FIRSTPRIVATE_INT on the libgomp oacc-*
side and just use the OpenMP firstprivate handling in omp-low.c?
Jakub
^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: OpenACC Firstprivate
2015-11-09 13:46 ` Jakub Jelinek
@ 2015-11-09 13:59 ` Nathan Sidwell
2015-11-09 14:06 ` Nathan Sidwell
2015-11-09 14:10 ` Jakub Jelinek
0 siblings, 2 replies; 11+ messages in thread
From: Nathan Sidwell @ 2015-11-09 13:59 UTC (permalink / raw)
To: Jakub Jelinek; +Cc: GCC Patches, Cesar Philippidis
On 11/09/15 08:46, Jakub Jelinek wrote:
> On Sat, Nov 07, 2015 at 08:50:28AM -0500, Nathan Sidwell wrote:
>> Index: gcc/gimplify.c
>> ===================================================================
>
> If you want to switch to hexadecimal, you should change all values
> in the enum to hexadecimal for consistency.
ok.
>>
>> /* Gimplify hashtable helper. */
>> @@ -377,6 +383,12 @@ new_omp_context (enum omp_region_type re
>> else
>> c->default_kind = OMP_CLAUSE_DEFAULT_UNSPECIFIED;
>>
>> + c->combined_loop = false;
>> + c->distribute = false;
>> + c->target_map_scalars_firstprivate = false;
>> + c->target_map_pointers_as_0len_arrays = false;
>> + c->target_firstprivatize_array_bases = false;
>
> Why this? c is XCNEW allocated, so zero initialized.
I presumed it necessary, as it was on the branch. will remove.
>
>> @@ -5667,11 +5682,13 @@ omp_add_variable (struct gimplify_omp_ct
>> /* We shouldn't be re-adding the decl with the same data
>> sharing class. */
>> gcc_assert ((n->value & GOVD_DATA_SHARE_CLASS & flags) == 0);
>> - /* The only combination of data sharing classes we should see is
>> - FIRSTPRIVATE and LASTPRIVATE. */
>> nflags = n->value | flags;
>> - gcc_assert ((nflags & GOVD_DATA_SHARE_CLASS)
>> - == (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE)
>> + /* The only combination of data sharing classes we should see is
>> + FIRSTPRIVATE and LASTPRIVATE. However, OpenACC permits
>> + reduction variables to be used in data sharing clauses. */
>> + gcc_assert ((ctx->region_type & ORT_ACC) != 0
>> + || ((nflags & GOVD_DATA_SHARE_CLASS)
>> + == (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE))
>> || (flags & GOVD_DATA_SHARE_CLASS) == 0);
>
> Are you sure you want to give up on any kind of consistency checks for
> OpenACC? If only reduction is special on OpenACC, perhaps you could tweak
> the assert for that instead? Something that can be done incrementally of
> course.
Will investigate (later)
>
>> +
>> + /* OpenMP doesn't look in outer contexts to find an
>> + enclosing data clause. */
>
> I'm puzzled by the comment. OpenMP does look in outer context for clauses
> that need that (pretty much all closes but private), that is the do_outer:
> recursion in omp_notice_variable. Say for firstprivate in order to copy (or
> copy construct) the private variable one needs the access to the outer
> context's var etc.).
> So perhaps it would help to document what you are doing here for OpenACC and
> why.
Ok. It seemed (and it may become clearer with default handling added), that
OpenACC and OpenMP scanned scopes in opposite orders. I remember trying to
get the ACC code to scan in the same order, but came up blank. Anyway, you're
right, it should say what OpenACC is trying.
> The main issue I have is with the omp-low.c changes.
> I see:
> "2.5.9
> private clause
> The private clause is allowed on the parallel construct; it declares that a copy of each
> item on the list will be created for each parallel gang.
>
> 2.5.10
> firstprivate clause
> The firstprivate clause is allowed on the parallel construct; it declares that a copy
> of each item on the list will be created for each parallel gang, and that the copy will be
> initialized with the value of that item on the host when the parallel construct is
> encountered."
>
> but looking at what you actually emit looks like standard present_copyin
> clause I think with a private variable defined in the region where the
> value of the present_copyin mapped variable is assigned to the private one.
> This I'm afraid performs often two copies rather than just one (one to copy
> the host value to the present_copyin mapped value, another one in the
> region),
I don't think that can be avoided. The host doesn't have control over when the
CTAs (a gang) start -- they may even be serialized onto the same physical HW.
So each gang has to initialize its own instance. Or did you mean something else?
> but more importantly, if the var is already mapped, you could
> initialize the private var with old data.
> Say
> int arr[64];
> // initialize arr
> #pragma acc data copyin (arr)
> {
> // modify arr on the host
> # pragma acc parallel firstprivate (arr)
> {
> ...
> }
> }
Hm, I suspect that is either ill formed or the std does not contemplate.
> Is that really what you want? If not, any reason not to implement
> GOMP_MAP_FIRSTPRIVATE and GOMP_MAP_FIRSTPRIVATE_INT on the libgomp oacc-*
> side and just use the OpenMP firstprivate handling in omp-low.c?
I would have to investigate ...
nathan
^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: OpenACC Firstprivate
2015-11-09 13:59 ` Nathan Sidwell
@ 2015-11-09 14:06 ` Nathan Sidwell
2015-11-09 14:10 ` Jakub Jelinek
1 sibling, 0 replies; 11+ messages in thread
From: Nathan Sidwell @ 2015-11-09 14:06 UTC (permalink / raw)
To: Jakub Jelinek; +Cc: GCC Patches, Cesar Philippidis
On 11/09/15 08:59, Nathan Sidwell wrote:
> On 11/09/15 08:46, Jakub Jelinek wrote:
>> On Sat, Nov 07, 2015 at 08:50:28AM -0500, Nathan Sidwell wrote:
>
>> Say
>> int arr[64];
>> // initialize arr
>> #pragma acc data copyin (arr)
>> {
>> // modify arr on the host
>> # pragma acc parallel firstprivate (arr)
>> {
>> ...
>> }
>> }
>
> Hm, I suspect that is either ill formed or the std does not contemplate.
just realized, there are two ways to consider the above.
1) it's ill formed. Once you've transferred data to the device, modifying it
on the host is unspecified. I'm having trouble finding words in the std that
actually say that though :(
2) on a system with shared physical global memory, the host modification would
be visiable on the device (possibly at an arbitrary point due to lack of
synchronization primitive?)
I don't think this changes 'why not use OpenMP's ...' question, because IIUC you
think that can be made to DTRT anyway?
nathan
^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: OpenACC Firstprivate
2015-11-09 13:59 ` Nathan Sidwell
2015-11-09 14:06 ` Nathan Sidwell
@ 2015-11-09 14:10 ` Jakub Jelinek
2015-11-09 14:46 ` Nathan Sidwell
1 sibling, 1 reply; 11+ messages in thread
From: Jakub Jelinek @ 2015-11-09 14:10 UTC (permalink / raw)
To: Nathan Sidwell; +Cc: GCC Patches, Cesar Philippidis
On Mon, Nov 09, 2015 at 08:59:15AM -0500, Nathan Sidwell wrote:
> >This I'm afraid performs often two copies rather than just one (one to copy
> >the host value to the present_copyin mapped value, another one in the
> >region),
>
> I don't think that can be avoided. The host doesn't have control over when
> the CTAs (a gang) start -- they may even be serialized onto the same
> physical HW. So each gang has to initialize its own instance. Or did you
> mean something else?
So, what is the scope of the private and firstprivate vars in OpenACC?
In OpenMP if a variable is private or firstprivate on the target construct,
unless further privatized in inner constructs it is really shared among all
the threads in all the teams (ro one var per all CTAs/workers in PTX terms).
Is that the case for OpenACC too, or are the vars e.g. private to each CTA
already or to each thread in each CTA, something different?
If they are shared by all CTAs, then you should hopefully be able to use the
GOMP_MAP_FIRSTPRIVATE{,_INT}, if not, then I'd say you should at least use
those to provide you the initializer data to initialize your private vars
from as a cheaper alternative to mapping.
Jakub
^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: OpenACC Firstprivate
2015-11-09 14:10 ` Jakub Jelinek
@ 2015-11-09 14:46 ` Nathan Sidwell
2015-11-10 14:13 ` Nathan Sidwell
0 siblings, 1 reply; 11+ messages in thread
From: Nathan Sidwell @ 2015-11-09 14:46 UTC (permalink / raw)
To: Jakub Jelinek; +Cc: GCC Patches, Cesar Philippidis
On 11/09/15 09:10, Jakub Jelinek wrote:
> On Mon, Nov 09, 2015 at 08:59:15AM -0500, Nathan Sidwell wrote:
>>> This I'm afraid performs often two copies rather than just one (one to copy
>>> the host value to the present_copyin mapped value, another one in the
>>> region),
>>
>> I don't think that can be avoided. The host doesn't have control over when
>> the CTAs (a gang) start -- they may even be serialized onto the same
>> physical HW. So each gang has to initialize its own instance. Or did you
>> mean something else?
>
> So, what is the scope of the private and firstprivate vars in OpenACC?
> In OpenMP if a variable is private or firstprivate on the target construct,
> unless further privatized in inner constructs it is really shared among all
> the threads in all the teams (ro one var per all CTAs/workers in PTX terms).
> Is that the case for OpenACC too, or are the vars e.g. private to each CTA
> already or to each thread in each CTA, something different?
> If they are shared by all CTAs, then you should hopefully be able to use the
> GOMP_MAP_FIRSTPRIVATE{,_INT}, if not, then I'd say you should at least use
> those to provide you the initializer data to initialize your private vars
> from as a cheaper alternative to mapping.
I'm going to try and get clarification, but I think the intent is to initialize
with the value seen on the device. Consider:
int foo = 0;
#pragma acc data copyin(foo)
{
#pragma acc parallel present(foo)
{
foo = 2;
}
if (expr){
#pragma update host (foo)
}
#pragma acc parallel firstprivate (foo)
{
// which initialization value?
}
}
Here we copy data to the device, then set it a distinct value there. We
conditionally update the host's instance from the device.
My thinking is that the intent of the firstprivate is to initialize with the
value known on the device (and behave as-if copyin, if it's not there). Not the
value most recently seen on the host -- the update clause could change that, and
may well be being used as a debugging aide, so it seems bizarre that it can
change program semantics in such a way.
^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: OpenACC Firstprivate
2015-11-09 14:46 ` Nathan Sidwell
@ 2015-11-10 14:13 ` Nathan Sidwell
2015-11-11 8:05 ` Jakub Jelinek
0 siblings, 1 reply; 11+ messages in thread
From: Nathan Sidwell @ 2015-11-10 14:13 UTC (permalink / raw)
To: Jakub Jelinek; +Cc: GCC Patches, Cesar Philippidis
[-- Attachment #1: Type: text/plain, Size: 1013 bytes --]
On 11/09/15 09:46, Nathan Sidwell wrote:
> I'm going to try and get clarification, but I think the intent is to initialize
> with the value seen on the device. Consider:
> My thinking is that the intent of the firstprivate is to initialize with the
> value known on the device (and behave as-if copyin, if it's not there). Not the
> value most recently seen on the host -- the update clause could change that, and
> may well be being used as a debugging aide, so it seems bizarre that it can
> change program semantics in such a way.
We believe my example is well formed. The data clauses transfer liveness of the
data from host to device (and vice versa). It is ill formed to manipulate the
data on the non-live system. firstprivate's intial value is taken from the
(statically determined) live location.
Unless I'm misunderstanding something about GOMP_MAP_FIRSTPRIVATE, using regular
target mapping is the right thing.
Here's an updated patch with the other two issues you noted fixed.
nathan
[-- Attachment #2: trunk-firstprivate-1110.patch --]
[-- Type: text/x-patch, Size: 21418 bytes --]
2015-11-10 Nathan Sidwell <nathan@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
gcc/
* gcc/gimplify.c (enum omp_region_type): Add ORT_ACC,
ORT_ACC_DATA, ORT_ACC_PARALLEL, ORT_ACC_KERNELS. Adjust ORT_NONE.
(gimple_add_tmp_var): Add ORT_ACC checks.
(gimplify_var_or_parm_decl): Likewise.
(omp_firstprivatize_variable): Likewise. Use ORT_TARGET_DATA as a
mask.
(omp_add_variable): Look in outer contexts for openacc and allow
reductions with other sharing. Add ORT_ACC and ORT_TARGET_DATA
checks.
(omp_notice_variable, omp_is_private, omp_check_private): Add
ORT_ACC checks.
(gimplify_scan_omp_clauses: Treat ORT_ACC as ORT_WORKSHARE.
Permit private openacc reductions.
(gimplify_oacc_cache): Specify ORT_ACC.
(gimplify_omp_workshare): Adjust OpenACC region types.
(gimplify_omp_target_update): Likewise.
* gcc/omp-low.c (scan_sharing_clauses): Remove Openacc
firstprivate sorry.
(lower-rec_input_clauses): Don't handle openacc firstprivate
references here.
(lower_omp_target): Emit initializers for openacc firstprivate vars.
gcc/testsuite/
* gfortran.dg/goacc/private-3.f95: Remove xfail.
* gfortran.dg/goacc/combined_loop.f90: Remove xfail.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c: Remove xfail.
* testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c: Remove xfail.
* testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c: New.
Index: gcc/gimplify.c
===================================================================
--- gcc/gimplify.c (revision 230107)
+++ gcc/gimplify.c (working copy)
@@ -95,22 +95,34 @@ enum gimplify_omp_var_data
enum omp_region_type
{
- ORT_WORKSHARE = 0,
- ORT_SIMD = 1,
- ORT_PARALLEL = 2,
- ORT_COMBINED_PARALLEL = 3,
- ORT_TASK = 4,
- ORT_UNTIED_TASK = 5,
- ORT_TEAMS = 8,
- ORT_COMBINED_TEAMS = 9,
+ ORT_WORKSHARE = 0x00,
+ ORT_SIMD = 0x01,
+
+ ORT_PARALLEL = 0x02,
+ ORT_COMBINED_PARALLEL = 0x03,
+
+ ORT_TASK = 0x04,
+ ORT_UNTIED_TASK = 0x05,
+
+ ORT_TEAMS = 0x08,
+ ORT_COMBINED_TEAMS = 0x09,
+
/* Data region. */
- ORT_TARGET_DATA = 16,
+ ORT_TARGET_DATA = 0x10,
+
/* Data region with offloading. */
- ORT_TARGET = 32,
- ORT_COMBINED_TARGET = 33,
+ ORT_TARGET = 0x20,
+ ORT_COMBINED_TARGET = 0x21,
+
+ /* OpenACC variants. */
+ ORT_ACC = 0x40, /* A generic OpenACC region. */
+ ORT_ACC_DATA = ORT_ACC | ORT_TARGET_DATA, /* Data construct. */
+ ORT_ACC_PARALLEL = ORT_ACC | ORT_TARGET, /* Parallel construct */
+ ORT_ACC_KERNELS = ORT_ACC | ORT_TARGET | 0x80, /* Kernels construct. */
+
/* Dummy OpenMP region, used to disable expansion of
DECL_VALUE_EXPRs in taskloop pre body. */
- ORT_NONE = 64
+ ORT_NONE = 0x100
};
/* Gimplify hashtable helper. */
@@ -689,7 +701,8 @@ gimple_add_tmp_var (tree tmp)
struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
while (ctx
&& (ctx->region_type == ORT_WORKSHARE
- || ctx->region_type == ORT_SIMD))
+ || ctx->region_type == ORT_SIMD
+ || ctx->region_type == ORT_ACC))
ctx = ctx->outer_context;
if (ctx)
omp_add_variable (ctx, tmp, GOVD_LOCAL | GOVD_SEEN);
@@ -1804,7 +1817,8 @@ gimplify_var_or_parm_decl (tree *expr_p)
struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
while (ctx
&& (ctx->region_type == ORT_WORKSHARE
- || ctx->region_type == ORT_SIMD))
+ || ctx->region_type == ORT_SIMD
+ || ctx->region_type == ORT_ACC))
ctx = ctx->outer_context;
if (!ctx && !nonlocal_vlas->add (decl))
{
@@ -5579,7 +5593,8 @@ omp_firstprivatize_variable (struct gimp
}
else if (ctx->region_type != ORT_WORKSHARE
&& ctx->region_type != ORT_SIMD
- && ctx->region_type != ORT_TARGET_DATA)
+ && ctx->region_type != ORT_ACC
+ && !(ctx->region_type & ORT_TARGET_DATA))
omp_add_variable (ctx, decl, GOVD_FIRSTPRIVATE);
ctx = ctx->outer_context;
@@ -5667,11 +5682,13 @@ omp_add_variable (struct gimplify_omp_ct
/* We shouldn't be re-adding the decl with the same data
sharing class. */
gcc_assert ((n->value & GOVD_DATA_SHARE_CLASS & flags) == 0);
- /* The only combination of data sharing classes we should see is
- FIRSTPRIVATE and LASTPRIVATE. */
nflags = n->value | flags;
- gcc_assert ((nflags & GOVD_DATA_SHARE_CLASS)
- == (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE)
+ /* The only combination of data sharing classes we should see is
+ FIRSTPRIVATE and LASTPRIVATE. However, OpenACC permits
+ reduction variables to be used in data sharing clauses. */
+ gcc_assert ((ctx->region_type & ORT_ACC) != 0
+ || ((nflags & GOVD_DATA_SHARE_CLASS)
+ == (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE))
|| (flags & GOVD_DATA_SHARE_CLASS) == 0);
n->value = nflags;
return;
@@ -5968,20 +5985,47 @@ omp_notice_variable (struct gimplify_omp
else if (is_scalar)
nflags |= GOVD_FIRSTPRIVATE;
}
- tree type = TREE_TYPE (decl);
- if (nflags == flags
- && gimplify_omp_ctxp->target_firstprivatize_array_bases
- && lang_hooks.decls.omp_privatize_by_reference (decl))
- type = TREE_TYPE (type);
- if (nflags == flags
- && !lang_hooks.types.omp_mappable_type (type))
- {
- error ("%qD referenced in target region does not have "
- "a mappable type", decl);
- nflags |= GOVD_MAP | GOVD_EXPLICIT;
+
+ struct gimplify_omp_ctx *octx = ctx->outer_context;
+ if ((ctx->region_type & ORT_ACC) && octx)
+ {
+ /* Look in outer OpenACC contexts, to see if there's a
+ data attribute for this variable. */
+ omp_notice_variable (octx, decl, in_code);
+
+ for (; octx; octx = octx->outer_context)
+ {
+ if (!(octx->region_type & (ORT_TARGET_DATA | ORT_TARGET)))
+ break;
+ splay_tree_node n2
+ = splay_tree_lookup (octx->variables,
+ (splay_tree_key) decl);
+ if (n2)
+ {
+ nflags |= GOVD_MAP;
+ goto found_outer;
+ }
+ }
}
- else if (nflags == flags)
- nflags |= GOVD_MAP;
+
+ {
+ tree type = TREE_TYPE (decl);
+
+ if (nflags == flags
+ && gimplify_omp_ctxp->target_firstprivatize_array_bases
+ && lang_hooks.decls.omp_privatize_by_reference (decl))
+ type = TREE_TYPE (type);
+ if (nflags == flags
+ && !lang_hooks.types.omp_mappable_type (type))
+ {
+ error ("%qD referenced in target region does not have "
+ "a mappable type", decl);
+ nflags |= GOVD_MAP | GOVD_EXPLICIT;
+ }
+ else if (nflags == flags)
+ nflags |= GOVD_MAP;
+ }
+ found_outer:
omp_add_variable (ctx, decl, nflags);
}
else
@@ -5998,7 +6042,8 @@ omp_notice_variable (struct gimplify_omp
{
if (ctx->region_type == ORT_WORKSHARE
|| ctx->region_type == ORT_SIMD
- || ctx->region_type == ORT_TARGET_DATA)
+ || ctx->region_type == ORT_ACC
+ || (ctx->region_type & ORT_TARGET_DATA) != 0)
goto do_outer;
flags = omp_default_clause (ctx, decl, in_code, flags);
@@ -6112,7 +6157,8 @@ omp_is_private (struct gimplify_omp_ctx
}
if (ctx->region_type != ORT_WORKSHARE
- && ctx->region_type != ORT_SIMD)
+ && ctx->region_type != ORT_SIMD
+ && ctx->region_type != ORT_ACC)
return false;
else if (ctx->outer_context)
return omp_is_private (ctx->outer_context, decl, simd);
@@ -6168,7 +6214,8 @@ omp_check_private (struct gimplify_omp_c
}
}
while (ctx->region_type == ORT_WORKSHARE
- || ctx->region_type == ORT_SIMD);
+ || ctx->region_type == ORT_SIMD
+ || ctx->region_type == ORT_ACC);
return false;
}
@@ -6311,7 +6358,8 @@ gimplify_scan_omp_clauses (tree *list_p,
omp_notice_variable (outer_ctx->outer_context, decl, true);
}
else if (outer_ctx
- && outer_ctx->region_type == ORT_WORKSHARE
+ && (outer_ctx->region_type == ORT_WORKSHARE
+ || outer_ctx->region_type == ORT_ACC)
&& outer_ctx->combined_loop
&& splay_tree_lookup (outer_ctx->variables,
(splay_tree_key) decl) == NULL
@@ -6335,7 +6383,9 @@ gimplify_scan_omp_clauses (tree *list_p,
goto do_add;
case OMP_CLAUSE_REDUCTION:
flags = GOVD_REDUCTION | GOVD_SEEN | GOVD_EXPLICIT;
- check_non_private = "reduction";
+ /* OpenACC permits reductions on private variables. */
+ if (!(region_type & ORT_ACC))
+ check_non_private = "reduction";
decl = OMP_CLAUSE_DECL (c);
if (TREE_CODE (decl) == MEM_REF)
{
@@ -7704,7 +7754,7 @@ gimplify_oacc_cache (tree *expr_p, gimpl
{
tree expr = *expr_p;
- gimplify_scan_omp_clauses (&OACC_CACHE_CLAUSES (expr), pre_p, ORT_WORKSHARE,
+ gimplify_scan_omp_clauses (&OACC_CACHE_CLAUSES (expr), pre_p, ORT_ACC,
OACC_CACHE);
gimplify_adjust_omp_clauses (pre_p, &OACC_CACHE_CLAUSES (expr), OACC_CACHE);
@@ -7833,7 +7883,9 @@ gimplify_omp_for (tree *expr_p, gimple_s
case OMP_FOR:
case CILK_FOR:
case OMP_DISTRIBUTE:
+ break;
case OACC_LOOP:
+ ort = ORT_ACC;
break;
case OMP_TASKLOOP:
if (find_omp_clause (OMP_FOR_CLAUSES (for_stmt), OMP_CLAUSE_UNTIED))
@@ -8895,10 +8947,14 @@ gimplify_omp_workshare (tree *expr_p, gi
ort = OMP_TARGET_COMBINED (expr) ? ORT_COMBINED_TARGET : ORT_TARGET;
break;
case OACC_KERNELS:
+ ort = ORT_ACC_KERNELS;
+ break;
case OACC_PARALLEL:
- ort = ORT_TARGET;
+ ort = ORT_ACC_PARALLEL;
break;
case OACC_DATA:
+ ort = ORT_ACC_DATA;
+ break;
case OMP_TARGET_DATA:
ort = ORT_TARGET_DATA;
break;
@@ -8920,7 +8976,7 @@ gimplify_omp_workshare (tree *expr_p, gi
pop_gimplify_context (g);
else
pop_gimplify_context (NULL);
- if (ort == ORT_TARGET_DATA)
+ if ((ort & ORT_TARGET_DATA) != 0)
{
enum built_in_function end_ix;
switch (TREE_CODE (expr))
@@ -8995,17 +9051,18 @@ gimplify_omp_target_update (tree *expr_p
tree expr = *expr_p;
int kind;
gomp_target *stmt;
+ enum omp_region_type ort = ORT_WORKSHARE;
switch (TREE_CODE (expr))
{
case OACC_ENTER_DATA:
- kind = GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA;
- break;
case OACC_EXIT_DATA:
kind = GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA;
+ ort = ORT_ACC;
break;
case OACC_UPDATE:
kind = GF_OMP_TARGET_KIND_OACC_UPDATE;
+ ort = ORT_ACC;
break;
case OMP_TARGET_UPDATE:
kind = GF_OMP_TARGET_KIND_UPDATE;
@@ -9020,7 +9077,7 @@ gimplify_omp_target_update (tree *expr_p
gcc_unreachable ();
}
gimplify_scan_omp_clauses (&OMP_STANDALONE_CLAUSES (expr), pre_p,
- ORT_WORKSHARE, TREE_CODE (expr));
+ ort, TREE_CODE (expr));
gimplify_adjust_omp_clauses (pre_p, &OMP_STANDALONE_CLAUSES (expr),
TREE_CODE (expr));
stmt = gimple_build_omp_target (NULL, kind, OMP_STANDALONE_CLAUSES (expr));
Index: gcc/omp-low.c
===================================================================
--- gcc/omp-low.c (revision 230107)
+++ gcc/omp-low.c (working copy)
@@ -1896,12 +1896,6 @@ scan_sharing_clauses (tree clauses, omp_
/* FALLTHRU */
case OMP_CLAUSE_FIRSTPRIVATE:
- if (is_gimple_omp_oacc (ctx->stmt))
- {
- sorry ("clause not supported yet");
- break;
- }
- /* FALLTHRU */
case OMP_CLAUSE_LINEAR:
decl = OMP_CLAUSE_DECL (c);
do_private:
@@ -2167,12 +2161,6 @@ scan_sharing_clauses (tree clauses, omp_
/* FALLTHRU */
case OMP_CLAUSE_FIRSTPRIVATE:
- if (is_gimple_omp_oacc (ctx->stmt))
- {
- sorry ("clause not supported yet");
- break;
- }
- /* FALLTHRU */
case OMP_CLAUSE_PRIVATE:
case OMP_CLAUSE_LINEAR:
case OMP_CLAUSE_IS_DEVICE_PTR:
@@ -4684,7 +4672,7 @@ lower_rec_input_clauses (tree clauses, g
gimplify_assign (ptr, x, ilist);
}
}
- else if (is_reference (var))
+ else if (is_reference (var) && !is_oacc_parallel (ctx))
{
/* For references that are being privatized for Fortran,
allocate new backing storage for the new pointer
@@ -14911,7 +14899,7 @@ lower_omp_target (gimple_stmt_iterator *
tree child_fn, t, c;
gomp_target *stmt = as_a <gomp_target *> (gsi_stmt (*gsi_p));
gbind *tgt_bind, *bind, *dep_bind = NULL;
- gimple_seq tgt_body, olist, ilist, new_body;
+ gimple_seq tgt_body, olist, ilist, fplist, new_body;
location_t loc = gimple_location (stmt);
bool offloaded, data_region;
unsigned int map_cnt = 0;
@@ -14963,6 +14951,7 @@ lower_omp_target (gimple_stmt_iterator *
child_fn = ctx->cb.dst_fn;
push_gimplify_context ();
+ fplist = NULL;
for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
switch (OMP_CLAUSE_CODE (c))
@@ -15007,6 +14996,7 @@ lower_omp_target (gimple_stmt_iterator *
/* FALLTHRU */
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
+ oacc_firstprivate:
var = OMP_CLAUSE_DECL (c);
if (!DECL_P (var))
{
@@ -15029,6 +15019,7 @@ lower_omp_target (gimple_stmt_iterator *
}
if (offloaded
+ && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
{
@@ -15057,17 +15048,40 @@ lower_omp_target (gimple_stmt_iterator *
x = build_receiver_ref (var, true, ctx);
tree new_var = lookup_decl (var, ctx);
- if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
&& TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
x = build_simple_mem_ref (x);
- SET_DECL_VALUE_EXPR (new_var, x);
- DECL_HAS_VALUE_EXPR_P (new_var) = 1;
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
+ {
+ gcc_assert (is_gimple_omp_oacc (ctx->stmt));
+ if (is_reference (new_var))
+ {
+ /* Create a local object to hold the instance
+ value. */
+ tree inst = create_tmp_var
+ (TREE_TYPE (TREE_TYPE (new_var)),
+ IDENTIFIER_POINTER (DECL_NAME (new_var)));
+ gimplify_assign (inst, fold_indirect_ref (x), &fplist);
+ x = build_fold_addr_expr (inst);
+ }
+ gimplify_assign (new_var, x, &fplist);
+ }
+ else if (DECL_P (new_var))
+ {
+ SET_DECL_VALUE_EXPR (new_var, x);
+ DECL_HAS_VALUE_EXPR_P (new_var) = 1;
+ }
+ else
+ gcc_unreachable ();
}
map_cnt++;
break;
case OMP_CLAUSE_FIRSTPRIVATE:
+ if (is_oacc_parallel (ctx))
+ goto oacc_firstprivate;
map_cnt++;
var = OMP_CLAUSE_DECL (c);
if (!is_reference (var)
@@ -15092,6 +15106,8 @@ lower_omp_target (gimple_stmt_iterator *
break;
case OMP_CLAUSE_PRIVATE:
+ if (is_gimple_omp_oacc (ctx->stmt))
+ break;
var = OMP_CLAUSE_DECL (c);
if (is_variable_sized (var))
{
@@ -15195,9 +15211,11 @@ lower_omp_target (gimple_stmt_iterator *
default:
break;
+
case OMP_CLAUSE_MAP:
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
+ oacc_firstprivate_map:
nc = c;
ovar = OMP_CLAUSE_DECL (c);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
@@ -15248,9 +15266,9 @@ lower_omp_target (gimple_stmt_iterator *
x = build_sender_ref (ovar, ctx);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
- && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
- && TREE_CODE (TREE_TYPE (ovar)) == ARRAY_TYPE)
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
+ && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
+ && TREE_CODE (TREE_TYPE (ovar)) == ARRAY_TYPE)
{
gcc_assert (offloaded);
tree avar
@@ -15261,6 +15279,15 @@ lower_omp_target (gimple_stmt_iterator *
avar = build_fold_addr_expr (avar);
gimplify_assign (x, avar, &ilist);
}
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
+ {
+ gcc_assert (is_gimple_omp_oacc (ctx->stmt));
+ if (!is_reference (var))
+ var = build_fold_addr_expr (var);
+ else
+ talign = TYPE_ALIGN_UNIT (TREE_TYPE (TREE_TYPE (ovar)));
+ gimplify_assign (x, var, &ilist);
+ }
else if (is_gimple_reg (var))
{
gcc_assert (offloaded);
@@ -15289,7 +15316,17 @@ lower_omp_target (gimple_stmt_iterator *
gimplify_assign (x, var, &ilist);
}
}
- s = OMP_CLAUSE_SIZE (c);
+ s = NULL_TREE;
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
+ {
+ gcc_checking_assert (is_gimple_omp_oacc (ctx->stmt));
+ s = TREE_TYPE (ovar);
+ if (TREE_CODE (s) == REFERENCE_TYPE)
+ s = TREE_TYPE (s);
+ s = TYPE_SIZE_UNIT (s);
+ }
+ else
+ s = OMP_CLAUSE_SIZE (c);
if (s == NULL_TREE)
s = TYPE_SIZE_UNIT (TREE_TYPE (ovar));
s = fold_convert (size_type_node, s);
@@ -15330,6 +15367,11 @@ lower_omp_target (gimple_stmt_iterator *
tkind_zero = tkind;
}
break;
+ case OMP_CLAUSE_FIRSTPRIVATE:
+ gcc_checking_assert (is_gimple_omp_oacc (ctx->stmt));
+ tkind = GOMP_MAP_TO;
+ tkind_zero = tkind;
+ break;
case OMP_CLAUSE_TO:
tkind = GOMP_MAP_TO;
tkind_zero = tkind;
@@ -15369,6 +15411,8 @@ lower_omp_target (gimple_stmt_iterator *
break;
case OMP_CLAUSE_FIRSTPRIVATE:
+ if (is_oacc_parallel (ctx))
+ goto oacc_firstprivate_map;
ovar = OMP_CLAUSE_DECL (c);
if (is_reference (ovar))
talign = TYPE_ALIGN_UNIT (TREE_TYPE (TREE_TYPE (ovar)));
@@ -15543,6 +15587,7 @@ lower_omp_target (gimple_stmt_iterator *
gimple_seq_add_stmt (&new_body,
gimple_build_assign (ctx->receiver_decl, t));
}
+ gimple_seq_add_seq (&new_body, fplist);
if (offloaded || data_region)
{
@@ -15554,6 +15599,8 @@ lower_omp_target (gimple_stmt_iterator *
default:
break;
case OMP_CLAUSE_FIRSTPRIVATE:
+ if (is_gimple_omp_oacc (ctx->stmt))
+ break;
var = OMP_CLAUSE_DECL (c);
if (is_reference (var)
|| is_gimple_reg_type (TREE_TYPE (var)))
@@ -15639,6 +15686,8 @@ lower_omp_target (gimple_stmt_iterator *
}
break;
case OMP_CLAUSE_PRIVATE:
+ if (is_gimple_omp_oacc (ctx->stmt))
+ break;
var = OMP_CLAUSE_DECL (c);
if (is_reference (var))
{
@@ -15727,7 +15776,7 @@ lower_omp_target (gimple_stmt_iterator *
/* Handle GOMP_MAP_FIRSTPRIVATE_{POINTER,REFERENCE} in second pass,
so that firstprivate vars holding OMP_CLAUSE_SIZE if needed
are already handled. */
- for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
+ for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
switch (OMP_CLAUSE_CODE (c))
{
tree var;
Index: gcc/testsuite/gfortran.dg/goacc/private-3.f95
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/private-3.f95 (revision 230107)
+++ gcc/testsuite/gfortran.dg/goacc/private-3.f95 (working copy)
@@ -1,6 +1,4 @@
! { dg-do compile }
-! <http://news.gmane.org/find-root.php?message_id=%3C563B78B5.5090506%40acm.org%3E>
-! { dg-xfail-if "TODO" { *-*-* } }
! test for private variables in a reduction clause
Index: gcc/testsuite/gfortran.dg/goacc/combined_loop.f90
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/combined_loop.f90 (revision 230107)
+++ gcc/testsuite/gfortran.dg/goacc/combined_loop.f90 (working copy)
@@ -1,6 +1,4 @@
! { dg-do compile }
-! <http://news.gmane.org/find-root.php?message_id=%3C563B78B5.5090506%40acm.org%3E>
-! { dg-xfail-if "TODO" { *-*-* } }
!
! PR fortran/64726
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c (revision 230107)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c (working copy)
@@ -1,7 +1,5 @@
/* { dg-do run } */
/* { dg-additional-options "-O2" */
-/* <http://news.gmane.org/find-root.php?message_id=%3C563B78B5.5090506%40acm.org%3E>
- { dg-xfail-if "TODO" { *-*-* } } */
#include <stdio.h>
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c (revision 230107)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c (working copy)
@@ -1,7 +1,5 @@
/* { dg-do run } */
/* { dg-additional-options "-O2" */
-/* <http://news.gmane.org/find-root.php?message_id=%3C563B78B5.5090506%40acm.org%3E>
- { dg-xfail-if "TODO" { *-*-* } } */
#include <stdio.h>
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c (revision 0)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c (working copy)
@@ -0,0 +1,41 @@
+/* { dg-do run } */
+
+#include <openacc.h>
+
+int main ()
+{
+ int ok = 1;
+ int val = 2;
+ int ary[32];
+ int ondev = 0;
+
+ for (int i = 0; i < 32; i++)
+ ary[i] = ~0;
+
+#pragma acc parallel num_gangs (32) copy (ok) firstprivate (val) copy(ary, ondev)
+ {
+ ondev = acc_on_device (acc_device_not_host);
+#pragma acc loop gang(static:1)
+ for (unsigned i = 0; i < 32; i++)
+ {
+ if (val != 2)
+ ok = 0;
+ val += i;
+ ary[i] = val;
+ }
+ }
+
+ if (ondev)
+ {
+ if (!ok)
+ return 1;
+ if (val != 2)
+ return 1;
+
+ for (int i = 0; i < 32; i++)
+ if (ary[i] != 2 + i)
+ return 1;
+ }
+
+ return 0;
+}
^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: OpenACC Firstprivate
2015-11-10 14:13 ` Nathan Sidwell
@ 2015-11-11 8:05 ` Jakub Jelinek
2015-11-11 13:44 ` Nathan Sidwell
2015-11-11 13:52 ` [gomp4] Rework gimplifyier region flags Nathan Sidwell
0 siblings, 2 replies; 11+ messages in thread
From: Jakub Jelinek @ 2015-11-11 8:05 UTC (permalink / raw)
To: Nathan Sidwell; +Cc: GCC Patches, Cesar Philippidis
On Tue, Nov 10, 2015 at 09:12:55AM -0500, Nathan Sidwell wrote:
> + /* Create a local object to hold the instance
> + value. */
> + tree inst = create_tmp_var
> + (TREE_TYPE (TREE_TYPE (new_var)),
> + IDENTIFIER_POINTER (DECL_NAME (new_var)));
Can you please rewrite this as:
tree type = TREE_TYPE (TREE_TYPE (new_var));
tree n = DECL_NAME (new_var);
tree inst = create_tmp_var (type, IDENTIFIER_POINTER (n));
or so (perhaps
const char *name
= IDENTIFIER_POINTER (DECL_NAME (new_var));
instead but then it takes one more line)?
I really don't like line breaks before opening ( unless really
necessary.
Otherwise LGTM.
Jakub
^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: OpenACC Firstprivate
2015-11-11 8:05 ` Jakub Jelinek
@ 2015-11-11 13:44 ` Nathan Sidwell
2015-11-12 9:00 ` Thomas Schwinge
2015-11-11 13:52 ` [gomp4] Rework gimplifyier region flags Nathan Sidwell
1 sibling, 1 reply; 11+ messages in thread
From: Nathan Sidwell @ 2015-11-11 13:44 UTC (permalink / raw)
To: Jakub Jelinek; +Cc: GCC Patches, Cesar Philippidis
On 11/11/15 03:04, Jakub Jelinek wrote:
> On Tue, Nov 10, 2015 at 09:12:55AM -0500, Nathan Sidwell wrote:
>> + /* Create a local object to hold the instance
>> + value. */
>> + tree inst = create_tmp_var
>> + (TREE_TYPE (TREE_TYPE (new_var)),
>> + IDENTIFIER_POINTER (DECL_NAME (new_var)));
>
> Can you please rewrite this as:
> tree type = TREE_TYPE (TREE_TYPE (new_var));
> tree n = DECL_NAME (new_var);
> tree inst = create_tmp_var (type, IDENTIFIER_POINTER (n));
> or so (perhaps
> const char *name
> = IDENTIFIER_POINTER (DECL_NAME (new_var));
> instead but then it takes one more line)?
> I really don't like line breaks before opening ( unless really
> necessary.
Oh, yeah you mentioned that before :)
>
> Otherwise LGTM.
thanks.
nathan
^ permalink raw reply [flat|nested] 11+ messages in thread
* [gomp4] Rework gimplifyier region flags
@ 2015-11-11 13:52 ` Nathan Sidwell
0 siblings, 0 replies; 11+ messages in thread
From: Nathan Sidwell @ 2015-11-11 13:52 UTC (permalink / raw)
To: GCC Patches
Cc: Thomas Schwinge, James Norris, Cesar Philippidis, Julian Brown
[-- Attachment #1: Type: text/plain, Size: 313 bytes --]
I've committed this patch to gomp4 to remove the openacc-specific enums from
gimplify_omp_ctx. Instead extending the existing omp_region_type enum. A
similar patch will shortly be applied to trunk, now Jakub.s approved it.
If you had patches relying on the old scheme, you'll need to update them.
nathan
[-- Attachment #2: gomp4-gimplify-1111.patch --]
[-- Type: text/x-patch, Size: 24677 bytes --]
2015-11-11 Nathan Sidwell <nathan@codesourcery.com>
* gimplify.c (enum gimplify_omp_var_data): Remove GOVD_FORCE_MAP.
(omp_region_type): Use hex. Add OpenACC members.
(omp_region_kind, acc_region_kind): Delete.
(gimplify_omp_ctx): Remove region_kind & acc_region_kind fields.
(new_omp_context): Adjust default_kind setting. Don't
reinitialize fiels.
(gimple_add_tmp_var): Add ORT_ACC check.
(gimplify_var_or_parm_decl): Likewise.
(omp_firstprivatize_variable): Likewise.
(omp_add_variable): Adjust OpenACC detection.
(oacc_default_clause): Reimplement.
(omp_notice_variable): Adjust OpenACC detection.
(gimplify_scan_omp_clauses): Remove region_kind arg. Adjust.
(gimplify_scan_omp_clause_1): Adjust OpenACC detection.
(gimmplify_oacc_cache, gimplify_oacc_declare,
gimplify_oacc_host_data, gimplify_omp_parallel): Adjust.
(gimplify_omp_for, gimplify_omp_workshare,
gimplify_omp_target_update): Adjust for OpenACC ORT flags.
(gimplify_expr): Likewise.
(gimplify_body): Simplify OpenACC declare handling.
Index: gimplify.c
===================================================================
--- gimplify.c (revision 230160)
+++ gimplify.c (working copy)
@@ -89,10 +89,8 @@ enum gimplify_omp_var_data
GOVD_USE_DEVICE = 1 << 17,
- GOVD_FORCE_MAP = 1 << 18,
-
/* OpenACC deviceptr clause. */
- GOVD_USE_DEVPTR = 1 << 19,
+ GOVD_USE_DEVPTR = 1 << 18,
GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
| GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
@@ -102,40 +100,37 @@ enum gimplify_omp_var_data
enum omp_region_type
{
- ORT_WORKSHARE = 0,
- ORT_SIMD = 1,
- ORT_PARALLEL = 2,
- ORT_COMBINED_PARALLEL = 3,
- ORT_TASK = 4,
- ORT_UNTIED_TASK = 5,
- ORT_TEAMS = 8,
- ORT_COMBINED_TEAMS = 9,
+ ORT_WORKSHARE = 0x00,
+ ORT_SIMD = 0x01,
+
+ ORT_PARALLEL = 0x02,
+ ORT_COMBINED_PARALLEL = 0x03,
+
+ ORT_TASK = 0x04,
+ ORT_UNTIED_TASK = 0x05,
+
+ ORT_TEAMS = 0x08,
+ ORT_COMBINED_TEAMS = 0x09,
+
/* Data region. */
- ORT_TARGET_DATA = 16,
+ ORT_TARGET_DATA = 0x10,
+
/* Data region with offloading. */
- ORT_TARGET = 32,
- ORT_COMBINED_TARGET = 33,
- /* An OpenACC host-data region. */
- ORT_HOST_DATA = 64,
- /* Dummy OpenMP region, used to disable expansion of
- DECL_VALUE_EXPRs in taskloop pre body. */
- ORT_NONE = 128
-};
+ ORT_TARGET = 0x20,
+ ORT_COMBINED_TARGET = 0x21,
-enum omp_region_kind
-{
- ORK_OMP,
- ORK_OACC,
- ORK_UNKNOWN
-};
+ ORT_HOST_DATA = 0x40,
-enum acc_region_kind
-{
- ARK_GENERAL, /* Default used for data, etc. regions. */
- ARK_PARALLEL, /* Parallel construct. */
- ARK_KERNELS, /* Kernels construct. */
- ARK_DECLARE, /* Declare directive. */
- ARK_UNKNOWN
+ /* OpenACC variants. */
+ ORT_ACC = 0x80, /* A generic OpenACC region. */
+ ORT_ACC_DATA = ORT_ACC | ORT_TARGET_DATA, /* Data construct. */
+ ORT_ACC_PARALLEL = ORT_ACC | ORT_TARGET, /* Parallel construct */
+ ORT_ACC_KERNELS = ORT_ACC | ORT_TARGET | 0x100, /* Kernels construct. */
+ ORT_ACC_HOST = ORT_ACC | ORT_HOST_DATA,
+
+ /* Dummy OpenMP region, used to disable expansion of
+ DECL_VALUE_EXPRs in taskloop pre body. */
+ ORT_NONE = 0x200
};
/* Gimplify hashtable helper. */
@@ -177,8 +172,6 @@ struct gimplify_omp_ctx
location_t location;
enum omp_clause_default_kind default_kind;
enum omp_region_type region_type;
- enum omp_region_kind region_kind;
- enum acc_region_kind acc_region_kind;
bool combined_loop;
bool distribute;
bool target_map_scalars_firstprivate;
@@ -404,19 +397,11 @@ new_omp_context (enum omp_region_type re
c->variables = splay_tree_new (splay_tree_compare_decl_uid, 0, 0);
c->privatized_types = new hash_set<tree>;
c->location = input_location;
- if ((region_type & (ORT_TASK | ORT_TARGET)) == 0)
+ c->region_type = region_type;
+ if ((region_type & ORT_TASK) == 0)
c->default_kind = OMP_CLAUSE_DEFAULT_SHARED;
else
c->default_kind = OMP_CLAUSE_DEFAULT_UNSPECIFIED;
- c->region_type = region_type;
- c->region_kind = ORK_UNKNOWN;
- c->acc_region_kind = ARK_UNKNOWN;
- c->combined_loop = false;
- c->distribute = false;
- c->target_map_scalars_firstprivate = false;
- c->target_map_pointers_as_0len_arrays = false;
- c->target_firstprivatize_array_bases = false;
- c->stmt = NULL;
return c;
}
@@ -730,7 +715,8 @@ gimple_add_tmp_var (tree tmp)
struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
while (ctx
&& (ctx->region_type == ORT_WORKSHARE
- || ctx->region_type == ORT_SIMD))
+ || ctx->region_type == ORT_SIMD
+ || ctx->region_type == ORT_ACC))
ctx = ctx->outer_context;
if (ctx)
omp_add_variable (ctx, tmp, GOVD_LOCAL | GOVD_SEEN);
@@ -1845,7 +1831,8 @@ gimplify_var_or_parm_decl (tree *expr_p)
struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
while (ctx
&& (ctx->region_type == ORT_WORKSHARE
- || ctx->region_type == ORT_SIMD))
+ || ctx->region_type == ORT_SIMD
+ || ctx->region_type == ORT_ACC))
ctx = ctx->outer_context;
if (!ctx && !nonlocal_vlas->add (decl))
{
@@ -5620,7 +5607,8 @@ omp_firstprivatize_variable (struct gimp
}
else if (ctx->region_type != ORT_WORKSHARE
&& ctx->region_type != ORT_SIMD
- && ctx->region_type != ORT_TARGET_DATA)
+ && ctx->region_type != ORT_ACC
+ && !(ctx->region_type & ORT_TARGET_DATA))
omp_add_variable (ctx, decl, GOVD_FIRSTPRIVATE);
ctx = ctx->outer_context;
@@ -5709,15 +5697,13 @@ omp_add_variable (struct gimplify_omp_ct
sharing class. */
gcc_assert ((n->value & GOVD_DATA_SHARE_CLASS & flags) == 0);
nflags = n->value | flags;
- if (ctx->region_kind != ORK_OACC)
- {
- /* The only combination of data sharing classes we should see is
- FIRSTPRIVATE and LASTPRIVATE. However, OpenACC permits
- reduction variables to be used in data sharing clauses. */
- gcc_assert ((nflags & GOVD_DATA_SHARE_CLASS)
- == (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE)
- || (flags & GOVD_DATA_SHARE_CLASS) == 0);
- }
+ /* The only combination of data sharing classes we should see is
+ FIRSTPRIVATE and LASTPRIVATE. However, OpenACC permits
+ reduction variables to be used in data sharing clauses. */
+ gcc_assert ((ctx->region_type & ORT_ACC) != 0
+ || ((nflags & GOVD_DATA_SHARE_CLASS)
+ == (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE))
+ || (flags & GOVD_DATA_SHARE_CLASS) == 0);
n->value = nflags;
return;
}
@@ -5954,54 +5940,51 @@ device_resident_p (tree decl)
static unsigned
oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
{
- switch (ctx->default_kind)
+ const char *rkind;
+ bool on_device = is_global_var (decl) && device_resident_p (decl);
+
+ if (on_device)
+ flags |= GOVD_MAP_TO_ONLY;
+
+ switch (ctx->region_type)
{
- default: gcc_unreachable ();
-
- case OMP_CLAUSE_DEFAULT_NONE:
- {
- const char *rkind;
+ default:
+ gcc_unreachable ();
- switch (ctx->acc_region_kind)
- {
- case ARK_PARALLEL: rkind = "parallel"; break;
- case ARK_KERNELS: rkind = "kernels"; break;
- default: gcc_unreachable ();
- }
- error ("%qE not specified in enclosing OpenACC %s construct",
- DECL_NAME (lang_hooks.decls.omp_report_decl (decl)), rkind);
- error_at (ctx->location, "enclosing OpenACC %s construct", rkind);
- }
- /* FALLTHRU. */
+ case ORT_ACC_KERNELS:
+ /* Everything under kernels are default 'present_or_copy'. */
+ flags |= GOVD_MAP;
+ rkind = "kernels";
+ break;
- case OMP_CLAUSE_DEFAULT_UNSPECIFIED:
+ case ORT_ACC_PARALLEL:
{
- if (is_global_var (decl) && device_resident_p (decl))
- flags |= GOVD_MAP_TO_ONLY | GOVD_MAP;
- else if (ctx->acc_region_kind == ARK_KERNELS)
- /* Everything under kernels are default 'copy'. */
- flags |= GOVD_FORCE_MAP | GOVD_MAP;
- else if (ctx->acc_region_kind == ARK_PARALLEL)
- {
- tree type = TREE_TYPE (decl);
+ tree type = TREE_TYPE (decl);
- if (TREE_CODE (type) == REFERENCE_TYPE
- || POINTER_TYPE_P (type))
- type = TREE_TYPE (type);
-
- if (AGGREGATE_TYPE_P (type))
- /* Aggregates default to 'present_or_copy'. */
- flags |= GOVD_MAP;
- else
- /* Scalars default to 'firstprivate'. */
- flags |= GOVD_FIRSTPRIVATE;
- }
+ if (TREE_CODE (type) == REFERENCE_TYPE
+ || POINTER_TYPE_P (type))
+ type = TREE_TYPE (type);
+
+ if (on_device || AGGREGATE_TYPE_P (type))
+ /* Aggregates default to 'present_or_copy'. */
+ flags |= GOVD_MAP;
else
- gcc_unreachable ();
+ /* Scalars default to 'firstprivate'. */
+ flags |= GOVD_FIRSTPRIVATE;
+ rkind = "parallel";
}
- break;
+ break;
}
-
+
+ if (ctx->default_kind == OMP_CLAUSE_DEFAULT_NONE)
+ {
+ error ("%qE not specified in enclosing OpenACC %s construct",
+ DECL_NAME (lang_hooks.decls.omp_report_decl (decl)), rkind);
+ error_at (ctx->location, "enclosing OpenACC %s construct", rkind);
+ }
+ else
+ gcc_checking_assert (ctx->default_kind == OMP_CLAUSE_DEFAULT_SHARED);
+
return flags;
}
@@ -6091,15 +6074,13 @@ omp_notice_variable (struct gimplify_omp
nflags |= GOVD_FIRSTPRIVATE;
}
- tree type = NULL_TREE;
-
- /* OpenMP doesn't look in outer contexts to find an
- enclosing data clause. */
struct gimplify_omp_ctx *octx = ctx->outer_context;
- if (ctx->region_kind == ORK_OACC && octx)
+ if ((ctx->region_type & ORT_ACC) && octx)
{
+ /* Look in outer OpenACC contexts, to see if there's a
+ data attribute for this variable. */
omp_notice_variable (octx, decl, in_code);
-
+
for (; octx; octx = octx->outer_context)
{
if (octx->region_type & ORT_HOST_DATA)
@@ -6117,27 +6098,30 @@ omp_notice_variable (struct gimplify_omp
}
}
- type = TREE_TYPE (decl);
- if (nflags == flags
- && gimplify_omp_ctxp->target_firstprivatize_array_bases
- && lang_hooks.decls.omp_privatize_by_reference (decl))
- type = TREE_TYPE (type);
- if (nflags == flags
- && !lang_hooks.types.omp_mappable_type (type,
- (ctx->region_kind
- == ORK_OACC)))
- {
- error ("%qD referenced in target region does not have "
- "a mappable type", decl);
- nflags |= GOVD_MAP | GOVD_EXPLICIT;
- }
- else if (nflags == flags)
- {
- if (ctx->region_kind == ORK_OACC)
- nflags = oacc_default_clause (ctx, decl, flags);
- else
- nflags |= GOVD_MAP;
- }
+ {
+ tree type = TREE_TYPE (decl);
+
+ if (nflags == flags
+ && gimplify_omp_ctxp->target_firstprivatize_array_bases
+ && lang_hooks.decls.omp_privatize_by_reference (decl))
+ type = TREE_TYPE (type);
+ if (nflags == flags
+ && !lang_hooks.types.omp_mappable_type (type,
+ (ctx->region_type
+ & ORT_ACC) != 0))
+ {
+ error ("%qD referenced in target region does not have "
+ "a mappable type", decl);
+ nflags |= GOVD_MAP | GOVD_EXPLICIT;
+ }
+ else if (nflags == flags)
+ {
+ if (ctx->region_type & ORT_ACC)
+ nflags = oacc_default_clause (ctx, decl, flags);
+ else
+ nflags |= GOVD_MAP;
+ }
+ }
found_outer:
omp_add_variable (ctx, decl, nflags);
}
@@ -6155,7 +6139,8 @@ omp_notice_variable (struct gimplify_omp
{
if (ctx->region_type == ORT_WORKSHARE
|| ctx->region_type == ORT_SIMD
- || ctx->region_type == ORT_TARGET_DATA)
+ || ctx->region_type == ORT_ACC
+ || (ctx->region_type & ORT_TARGET_DATA) != 0)
goto do_outer;
flags = omp_default_clause (ctx, decl, in_code, flags);
@@ -6269,7 +6254,8 @@ omp_is_private (struct gimplify_omp_ctx
}
if (ctx->region_type != ORT_WORKSHARE
- && ctx->region_type != ORT_SIMD)
+ && ctx->region_type != ORT_SIMD
+ && ctx->region_type != ORT_ACC)
return false;
else if (ctx->outer_context)
return omp_is_private (ctx->outer_context, decl, simd);
@@ -6325,7 +6311,8 @@ omp_check_private (struct gimplify_omp_c
}
}
while (ctx->region_type == ORT_WORKSHARE
- || ctx->region_type == ORT_SIMD);
+ || ctx->region_type == ORT_SIMD
+ || ctx->region_type == ORT_ACC);
return false;
}
@@ -6381,8 +6368,7 @@ find_decl_expr (tree *tp, int *walk_subt
static void
gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
enum omp_region_type region_type,
- enum tree_code code,
- enum omp_region_kind region_kind)
+ enum tree_code code)
{
struct gimplify_omp_ctx *ctx, *outer_ctx;
tree c, clauses = *list_p;
@@ -6393,7 +6379,6 @@ gimplify_scan_omp_clauses (tree *list_p,
ctx = new_omp_context (region_type);
outer_ctx = ctx->outer_context;
- ctx->region_kind = region_kind;
redvec.create (8);
if (code == OMP_TARGET && !lang_GNU_Fortran ())
{
@@ -6473,7 +6458,8 @@ gimplify_scan_omp_clauses (tree *list_p,
omp_notice_variable (outer_ctx->outer_context, decl, true);
}
else if (outer_ctx
- && outer_ctx->region_type == ORT_WORKSHARE
+ && (outer_ctx->region_type == ORT_WORKSHARE
+ || outer_ctx->region_type == ORT_ACC)
&& outer_ctx->combined_loop
&& splay_tree_lookup (outer_ctx->variables,
(splay_tree_key) decl) == NULL
@@ -6497,6 +6483,9 @@ gimplify_scan_omp_clauses (tree *list_p,
goto do_add;
case OMP_CLAUSE_REDUCTION:
flags = GOVD_REDUCTION | GOVD_SEEN | GOVD_EXPLICIT;
+ /* OpenACC permits reductions on private variables. */
+ if (!(region_type & ORT_ACC))
+ check_non_private = "reduction";
decl = OMP_CLAUSE_DECL (c);
if (TREE_CODE (decl) == MEM_REF)
{
@@ -6535,13 +6524,10 @@ gimplify_scan_omp_clauses (tree *list_p,
|| TREE_CODE (decl) == INDIRECT_REF)
decl = TREE_OPERAND (decl, 0);
}
- if ((region_kind == ORK_OACC) && ((region_type & ORT_TARGET) != 0)
+ if ((region_type & ORT_ACC) && ((region_type & ORT_TARGET) != 0)
&& (outer_ctx == NULL
- || (outer_ctx->region_kind == ORK_OACC
- && outer_ctx->region_type == ORT_TARGET_DATA)))
+ || outer_ctx->region_type == ORT_ACC_DATA))
redvec.safe_push (OMP_CLAUSE_DECL (c));
- if (region_kind != ORK_OACC)
- check_non_private = "reduction";
goto do_add_decl;
case OMP_CLAUSE_USE_DEVICE:
flags = GOVD_USE_DEVICE | GOVD_EXPLICIT;
@@ -7025,7 +7011,6 @@ gimplify_scan_omp_clauses (tree *list_p,
prev_list_p = list_p;
break;
}
-
flags = GOVD_MAP | GOVD_EXPLICIT;
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TO
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TOFROM)
@@ -7381,7 +7366,8 @@ gimplify_scan_omp_clauses (tree *list_p,
/* Add an implicit data-movement clause for an OpenACC parallel
reduction, if necessary. */
if (OMP_CLAUSE_CHAIN (c) == NULL && !processed_reductions
- && ((region_type & ORT_TARGET) != 0) && region_kind == ORK_OACC)
+ && ((region_type & ORT_TARGET) != 0)
+ && (region_type & ORT_ACC))
{
tree t;
@@ -7410,7 +7396,7 @@ gimplify_scan_omp_clauses (tree *list_p,
splay_tree_node n;
n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
- n->value |= GOVD_FORCE_MAP;
+ n->value |= GOVD_MAP;
}
}
@@ -7449,7 +7435,7 @@ gimplify_adjust_omp_clauses_1 (splay_tre
tree clause;
bool private_debug;
- if (flags & (GOVD_EXPLICIT | GOVD_LOCAL) && ((flags & GOVD_FORCE_MAP) == 0))
+ if (flags & (GOVD_EXPLICIT | GOVD_LOCAL))
return 0;
if ((flags & GOVD_SEEN) == 0)
return 0;
@@ -7466,7 +7452,7 @@ gimplify_adjust_omp_clauses_1 (splay_tre
!!(flags & GOVD_SHARED));
if (private_debug)
code = OMP_CLAUSE_PRIVATE;
- else if (flags & (GOVD_MAP | GOVD_FORCE_MAP))
+ else if (flags & GOVD_MAP)
code = OMP_CLAUSE_MAP;
else if (flags & GOVD_SHARED)
{
@@ -7533,11 +7519,9 @@ gimplify_adjust_omp_clauses_1 (splay_tre
else if (code == OMP_CLAUSE_MAP)
{
OMP_CLAUSE_SET_MAP_KIND (clause,
- flags & GOVD_MAP_TO_ONLY ? GOMP_MAP_TO
- : (flags & GOVD_FORCE_MAP
- ? GOMP_MAP_FORCE_TOFROM
- : GOMP_MAP_TOFROM));
-
+ flags & GOVD_MAP_TO_ONLY
+ ? GOMP_MAP_TO
+ : GOMP_MAP_TOFROM);
if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
{
@@ -7931,8 +7915,8 @@ gimplify_oacc_cache (tree *expr_p, gimpl
{
tree expr = *expr_p;
- gimplify_scan_omp_clauses (&OACC_CACHE_CLAUSES (expr), pre_p, ORT_WORKSHARE,
- OACC_CACHE, ORK_OACC);
+ gimplify_scan_omp_clauses (&OACC_CACHE_CLAUSES (expr), pre_p, ORT_ACC,
+ OACC_CACHE);
gimplify_adjust_omp_clauses (pre_p, &OACC_CACHE_CLAUSES (expr), OACC_CACHE);
/* TODO: Do something sensible with this information. */
@@ -7951,10 +7935,8 @@ gimplify_oacc_declare (tree *expr_p, gim
clauses = OACC_DECLARE_CLAUSES (expr);
- gimplify_scan_omp_clauses (&clauses, pre_p, ORT_TARGET_DATA, OACC_DECLARE,
- ORK_OACC);
+ gimplify_scan_omp_clauses (&clauses, pre_p, ORT_ACC_DATA, OACC_DECLARE);
- gimplify_omp_ctxp->acc_region_kind = ARK_DECLARE;
gimplify_omp_ctxp->stmt = NULL;
for (t = clauses; t; t = OMP_CLAUSE_CHAIN (t))
@@ -7982,8 +7964,7 @@ gimplify_oacc_declare (tree *expr_p, gim
{
struct gimplify_omp_ctx *c;
- gimplify_scan_omp_clauses (&clauses, pre_p, ORT_TARGET_DATA,
- OACC_DECLARE, ORK_OACC);
+ gimplify_scan_omp_clauses (&clauses, pre_p, ORT_ACC_DATA, OACC_DECLARE);
c = gimplify_omp_ctxp;
gimplify_omp_ctxp = c->outer_context;
@@ -8080,7 +8061,7 @@ gimplify_oacc_host_data (tree *expr_p, g
gimple_seq body = NULL;
gimplify_scan_omp_clauses (&OACC_HOST_DATA_CLAUSES (expr), pre_p,
- ORT_HOST_DATA, OACC_HOST_DATA, ORK_OACC);
+ ORT_ACC_HOST, OACC_HOST_DATA);
orig_body = OACC_HOST_DATA_BODY (expr);
@@ -8132,8 +8113,7 @@ gimplify_omp_parallel (tree *expr_p, gim
gimplify_scan_omp_clauses (&OMP_PARALLEL_CLAUSES (expr), pre_p,
OMP_PARALLEL_COMBINED (expr)
? ORT_COMBINED_PARALLEL
- : ORT_PARALLEL,
- OMP_PARALLEL, ORK_OMP);
+ : ORT_PARALLEL, OMP_PARALLEL);
push_gimplify_context ();
@@ -8170,8 +8150,7 @@ gimplify_omp_task (tree *expr_p, gimple_
gimplify_scan_omp_clauses (&OMP_TASK_CLAUSES (expr), pre_p,
find_omp_clause (OMP_TASK_CLAUSES (expr),
OMP_CLAUSE_UNTIED)
- ? ORT_UNTIED_TASK : ORT_TASK,
- OMP_TASK, ORK_OMP);
+ ? ORT_UNTIED_TASK : ORT_TASK, OMP_TASK);
push_gimplify_context ();
@@ -8312,7 +8291,6 @@ gimplify_omp_for (tree *expr_p, gimple_s
int i;
bitmap has_decl_expr = NULL;
enum omp_region_type ort = ORT_WORKSHARE;
- enum omp_region_kind ork;
orig_for_stmt = for_stmt = *expr_p;
@@ -8321,28 +8299,25 @@ gimplify_omp_for (tree *expr_p, gimple_s
case OMP_FOR:
case CILK_FOR:
case OMP_DISTRIBUTE:
- ork = ORK_OMP;
break;
case OACC_LOOP:
- ork = ORK_OACC;
+ ort = ORT_ACC;
break;
case OMP_TASKLOOP:
if (find_omp_clause (OMP_FOR_CLAUSES (for_stmt), OMP_CLAUSE_UNTIED))
ort = ORT_UNTIED_TASK;
else
ort = ORT_TASK;
- ork = ORK_OMP;
break;
case OMP_SIMD:
case CILK_SIMD:
ort = ORT_SIMD;
- ork = ORK_OMP;
break;
default:
gcc_unreachable ();
}
- if (ork == ORK_OACC)
+ if (ort == ORT_ACC)
localize_reductions (expr_p, false);
/* Set OMP_CLAUSE_LINEAR_NO_COPYIN flag on explicit linear
@@ -8376,7 +8351,7 @@ gimplify_omp_for (tree *expr_p, gimple_s
if (TREE_CODE (for_stmt) != OMP_TASKLOOP)
gimplify_scan_omp_clauses (&OMP_FOR_CLAUSES (for_stmt), pre_p, ort,
- TREE_CODE (for_stmt), ork);
+ TREE_CODE (for_stmt));
if (TREE_CODE (for_stmt) == OMP_DISTRIBUTE)
gimplify_omp_ctxp->distribute = true;
@@ -8484,7 +8459,7 @@ gimplify_omp_for (tree *expr_p, gimple_s
}
gimplify_scan_omp_clauses (&OMP_FOR_CLAUSES (orig_for_stmt), pre_p, ort,
- OMP_TASKLOOP, ork);
+ OMP_TASKLOOP);
}
if (orig_for_stmt != for_stmt)
@@ -9380,55 +9355,43 @@ gimplify_omp_workshare (tree *expr_p, gi
gimple *stmt;
gimple_seq body = NULL;
enum omp_region_type ort;
- enum omp_region_kind ork;
- enum acc_region_kind ark = ARK_GENERAL;
switch (TREE_CODE (expr))
{
case OMP_SECTIONS:
case OMP_SINGLE:
- ork = ORK_OMP;
ort = ORT_WORKSHARE;
break;
case OMP_TARGET:
- ork = ORK_OMP;
ort = OMP_TARGET_COMBINED (expr) ? ORT_COMBINED_TARGET : ORT_TARGET;
break;
case OACC_KERNELS:
- ark = ARK_KERNELS;
- ork = ORK_OACC;
- ort = ORT_TARGET;
+ ort = ORT_ACC_KERNELS;
break;
case OACC_PARALLEL:
- ark = ARK_PARALLEL;
- ork = ORK_OACC;
- ort = ORT_TARGET;
+ ort = ORT_ACC_PARALLEL;
break;
case OACC_DATA:
- ort = ORT_TARGET_DATA;
- ork = ORK_OACC;
+ ort = ORT_ACC_DATA;
break;
case OMP_TARGET_DATA:
- ork = ORK_OMP;
ort = ORT_TARGET_DATA;
break;
case OMP_TEAMS:
- ork = ORK_OMP;
ort = OMP_TEAMS_COMBINED (expr) ? ORT_COMBINED_TEAMS : ORT_TEAMS;
break;
default:
gcc_unreachable ();
}
gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ort,
- TREE_CODE (expr), ork);
+ TREE_CODE (expr));
if (TREE_CODE (expr) == OMP_TARGET)
optimize_target_teams (expr, pre_p);
if ((ort & (ORT_TARGET | ORT_TARGET_DATA)) != 0)
{
- gimplify_omp_ctxp->acc_region_kind = ark;
push_gimplify_context ();
- if (ork == ORK_OACC)
+ if (ort & ORT_ACC)
localize_reductions (expr_p, true);
gimple *g = gimplify_and_return_first (OMP_BODY (expr), &body);
@@ -9436,7 +9399,7 @@ gimplify_omp_workshare (tree *expr_p, gi
pop_gimplify_context (g);
else
pop_gimplify_context (NULL);
- if (ort == ORT_TARGET_DATA)
+ if ((ort & ORT_TARGET_DATA) != 0)
{
enum built_in_function end_ix;
switch (TREE_CODE (expr))
@@ -9511,39 +9474,33 @@ gimplify_omp_target_update (tree *expr_p
tree expr = *expr_p;
int kind;
gomp_target *stmt;
- enum omp_region_kind ork;
+ enum omp_region_type ort = ORT_WORKSHARE;
switch (TREE_CODE (expr))
{
case OACC_ENTER_DATA:
- kind = GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA;
- ork = ORK_OACC;
- break;
case OACC_EXIT_DATA:
kind = GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA;
- ork = ORK_OACC;
+ ort = ORT_ACC;
break;
case OACC_UPDATE:
kind = GF_OMP_TARGET_KIND_OACC_UPDATE;
- ork = ORK_OACC;
+ ort = ORT_ACC;
break;
case OMP_TARGET_UPDATE:
kind = GF_OMP_TARGET_KIND_UPDATE;
- ork = ORK_OMP;
break;
case OMP_TARGET_ENTER_DATA:
kind = GF_OMP_TARGET_KIND_ENTER_DATA;
- ork = ORK_OMP;
break;
case OMP_TARGET_EXIT_DATA:
kind = GF_OMP_TARGET_KIND_EXIT_DATA;
- ork = ORK_OMP;
break;
default:
gcc_unreachable ();
}
gimplify_scan_omp_clauses (&OMP_STANDALONE_CLAUSES (expr), pre_p,
- ORT_WORKSHARE, TREE_CODE (expr), ork);
+ ort, TREE_CODE (expr));
gimplify_adjust_omp_clauses (pre_p, &OMP_STANDALONE_CLAUSES (expr),
TREE_CODE (expr));
stmt = gimple_build_omp_target (NULL, kind, OMP_STANDALONE_CLAUSES (expr));
@@ -10655,8 +10612,7 @@ gimplify_expr (tree *expr_p, gimple_seq
break;
case OMP_CRITICAL:
gimplify_scan_omp_clauses (&OMP_CRITICAL_CLAUSES (*expr_p),
- pre_p, ORT_WORKSHARE, OMP_CRITICAL,
- ORK_OMP);
+ pre_p, ORT_WORKSHARE, OMP_CRITICAL);
gimplify_adjust_omp_clauses (pre_p,
&OMP_CRITICAL_CLAUSES (*expr_p),
OMP_CRITICAL);
@@ -11258,23 +11214,15 @@ gimplify_body (tree fndecl, bool do_parm
gimplify_seq_add_stmt (&seq, outer_stmt);
}
- if (flag_openacc && gimplify_omp_ctxp)
- {
- while (gimplify_omp_ctxp)
- {
- struct gimplify_omp_ctx *c;
-
- if (gimplify_omp_ctxp->acc_region_kind == ARK_DECLARE
- && gimplify_omp_ctxp->stmt)
- {
- gimplify_seq_add_stmt (&seq, gimplify_omp_ctxp->stmt);
- gimplify_omp_ctxp->stmt = NULL;
- }
+ if (flag_openacc)
+ while (gimplify_omp_ctxp)
+ {
+ if (gimplify_omp_ctxp->stmt)
+ gimplify_seq_add_stmt (&seq, gimplify_omp_ctxp->stmt);
- c = gimplify_omp_ctxp;
- gimplify_omp_ctxp = c->outer_context;
- delete_omp_context (c);
- }
+ struct gimplify_omp_ctx *c = gimplify_omp_ctxp;
+ gimplify_omp_ctxp = c->outer_context;
+ delete_omp_context (c);
}
/* The body must contain exactly one statement, a GIMPLE_BIND. If this is
^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: OpenACC Firstprivate
2015-11-11 13:44 ` Nathan Sidwell
@ 2015-11-12 9:00 ` Thomas Schwinge
0 siblings, 0 replies; 11+ messages in thread
From: Thomas Schwinge @ 2015-11-12 9:00 UTC (permalink / raw)
To: Nathan Sidwell; +Cc: GCC Patches, Cesar Philippidis, Jakub Jelinek
[-- Attachment #1: Type: text/plain, Size: 1602 bytes --]
Hi Nathan!
Merging back your trunk r230169 into gomp-4_0-branch, for the new
libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-2.c test, I'm
seeing the compiler diagnose as follows (compile with "-Wall -O2"):
source-gcc/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-2.c: In function 'main._omp_fn.1':
source-gcc/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-2.c:20:17: warning: 'val' is used uninitialized in this function [-Wuninitialized]
ok = val == 7;
^
source-gcc/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-2.c:9:7: note: 'val' was declared here
int val = 2;
^
..., and execution fails ("return 1" from main), so I XFAILed the
execution in the merge commit r230214 on gomp-4_0-branch. (..., and I
still think that it's a good idea to change the libgomp testsuite to run
with -Wall enabled...)
Do you have an idea what's going on? Given your preparatory "[gomp4]
Rework gimplifyier region flags",
<http://news.gmane.org/find-root.php?message_id=%3C56434833.7010703%40acm.org%3E>
(thanks!), the merge commit r230214 on gomp-4_0-branch didn't contain any
changes to gcc/gimplify.c, so that can't be it. It also can't be the
possibly inconsistent usage of gcc/omp-low.c:is_reference vs. "TREE_CODE
(TREE_TYPE ([...])) == REFERENCE_TYPE" in gcc/omp-low.c, because that
doesn't matter for C code anyway (no artificial REFERENCE_TYPEs
generated), right? So it must be some other change installed on
gomp-4_0-branch but not on trunk.
Grüße
Thomas
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]
^ permalink raw reply [flat|nested] 11+ messages in thread
end of thread, other threads:[~2015-11-12 9:00 UTC | newest]
Thread overview: 11+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2015-11-07 13:50 OpenACC Firstprivate Nathan Sidwell
2015-11-09 13:46 ` Jakub Jelinek
2015-11-09 13:59 ` Nathan Sidwell
2015-11-09 14:06 ` Nathan Sidwell
2015-11-09 14:10 ` Jakub Jelinek
2015-11-09 14:46 ` Nathan Sidwell
2015-11-10 14:13 ` Nathan Sidwell
2015-11-11 8:05 ` Jakub Jelinek
2015-11-11 13:44 ` Nathan Sidwell
2015-11-12 9:00 ` Thomas Schwinge
2015-11-11 13:52 ` [gomp4] Rework gimplifyier region flags Nathan Sidwell
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).