I've committed this patch to gomp4. The existing implementation of firstprivate
presumes the existence of memory at the CTA level. This patch does away with
that, treating firstprivate as thread-private variables initialized from the
host.
During development there was some fallout from declare handling, as that wasn't
creating the expected omp_region context object. The previous handling of
firstprivate just happened to work. Jim has been working on resolving that problem.
nathan
2015-08-03 Nathan Sidwell <nat...@codesourcery.com>
* gimplify.c (GOVD_GANGLOCAL): Delete.
(oacc_default_clause): Only derereference reference types. Mark
firstprivate as GOVD_FIRSTPRIVATE.
(gimplify_adjust_omp_clauses_1): Remove GANGLOCALL handling.
(gimplify_omp_for): Remove bogus OpenACC outer context lookup.
* omp-low.c (build_outer_var_ref): Simplify openacc outer ref
lookup.
(scan_sharing_clauses): Handle openacc firstprivate.
(lower_omp_target): Handle openacc firstprivate.
c/
* c-parser.c (c_parser_oacc_data_clause): Remove firstprivate
handling.
(c_parser_oac_all_clauses): Firstpribsste is a firstprivate
clause.
* c-typeck.c (c_finish_omp_clauses): Remove GANGLOCAL handling.
fortran/
* trans-openmp.c (gfc_trans_omp_clauses_1): Remove GANGLOCAL
handling.
* gfortran.h (OMP_MAP_GANGLOCAL): Delete.
(OMP_MAP_FORCE_TO_GANGLOCAL): Likewise.
* openmp.c (gfc_match_omp_clauses): Remove openacc specific
firstprivate handling.
testsuite/
* gfortran.dg/goacc/parallel-tree.f95: Remove ganglocal
expectation.
* gfortran.dg/goacc/list.f95: Stop expected firstprivate to be a
data clause.
* c-c++-common/goacc/firstprivate.c: Likewise.
cp/
* semantics.c (finish_omp_clauses): Remove OpenACC-specific
firstprivate handling.
* parser.c (cp_parser_oacc_data_clause): Remove firstprivate here.
(cp_parser_oacc_all_clauses): First private is a firstprivate clause.
Index: gcc/gimplify.c
===================================================================
--- gcc/gimplify.c (revision 226462)
+++ gcc/gimplify.c (working copy)
@@ -94,9 +94,6 @@ enum gimplify_omp_var_data
GOVD_FORCE_MAP = 1 << 16,
- /* Gang-local OpenACC variable. */
- GOVD_GANGLOCAL = 1 << 17,
-
/* OpenACC deviceptr clause. */
GOVD_USE_DEVPTR = 1 << 18,
@@ -5937,14 +5934,13 @@ oacc_default_clause (struct gimplify_omp
if (is_global_var (decl) && device_resident_p (decl))
flags |= GOVD_MAP_TO_ONLY | GOVD_MAP;
else if (ctx->acc_region_kind == ARK_KERNELS)
- /* Scalars under kernels are default 'copy'. */
+ /* 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);
- /* Should this be REFERENCE_TYPE_P? */
- if (POINTER_TYPE_P (type))
+ if (TREE_CODE (type) == REFERENCE_TYPE)
type = TREE_TYPE (type);
if (AGGREGATE_TYPE_P (type))
@@ -5952,12 +5948,12 @@ oacc_default_clause (struct gimplify_omp
flags |= GOVD_MAP;
else
/* Scalars default to 'firstprivate'. */
- flags |= GOVD_GANGLOCAL | GOVD_MAP_TO_ONLY | GOVD_MAP;
+ flags |= GOVD_FIRSTPRIVATE;
}
else
gcc_unreachable ();
}
- break;
+ break;
}
return flags;
@@ -6812,10 +6808,7 @@ gimplify_adjust_omp_clauses_1 (splay_tre
else if (code == OMP_CLAUSE_MAP)
{
OMP_CLAUSE_SET_MAP_KIND (clause,
- flags & GOVD_MAP_TO_ONLY
- ? (flags & GOVD_GANGLOCAL
- ? GOMP_MAP_FORCE_TO_GANGLOCAL
- : GOMP_MAP_TO)
+ flags & GOVD_MAP_TO_ONLY ? GOMP_MAP_TO
: (flags & GOVD_FORCE_MAP
? GOMP_MAP_FORCE_TOFROM
: GOMP_MAP_TOFROM));
@@ -7542,11 +7535,7 @@ gimplify_omp_for (tree *expr_p, gimple_s
else if (omp_is_private (gimplify_omp_ctxp, decl, 0))
omp_notice_variable (gimplify_omp_ctxp, decl, true);
else
- {
- if (ork == ORK_OACC && gimplify_omp_ctxp->outer_context)
- omp_notice_variable (gimplify_omp_ctxp->outer_context, decl, true);
- omp_add_variable (gimplify_omp_ctxp, decl, GOVD_PRIVATE | GOVD_SEEN);
- }
+ omp_add_variable (gimplify_omp_ctxp, decl, GOVD_PRIVATE | GOVD_SEEN);
/* If DECL is not a gimple register, create a temporary variable to act
as an iteration counter. This is valid, since DECL cannot be
Index: gcc/c/c-parser.c
===================================================================
--- gcc/c/c-parser.c (revision 226462)
+++ gcc/c/c-parser.c (working copy)
@@ -10719,9 +10719,6 @@ c_parser_oacc_data_clause (c_parser *par
case PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT:
kind = GOMP_MAP_DEVICE_RESIDENT;
break;
- case PRAGMA_OACC_CLAUSE_FIRSTPRIVATE:
- kind = GOMP_MAP_FORCE_TO_GANGLOCAL;
- break;
case PRAGMA_OACC_CLAUSE_HOST:
kind = GOMP_MAP_FORCE_FROM;
break;
@@ -12316,7 +12313,7 @@ c_parser_oacc_all_clauses (c_parser *par
c_name = "deviceptr";
break;
case PRAGMA_OACC_CLAUSE_FIRSTPRIVATE:
- clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+ clauses = c_parser_omp_clause_firstprivate (parser, clauses);
c_name = "firstprivate";
break;
case PRAGMA_OACC_CLAUSE_GANG:
Index: gcc/c/c-typeck.c
===================================================================
--- gcc/c/c-typeck.c (revision 226462)
+++ gcc/c/c-typeck.c (working copy)
@@ -12435,10 +12435,6 @@ c_finish_omp_clauses (tree clauses, bool
t = OMP_CLAUSE_DECL (c);
if (TREE_CODE (t) == TREE_LIST)
{
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_TO_GANGLOCAL))
- error_at (OMP_CLAUSE_LOCATION (c),
- "subarrays are not permitted in firstprivate");
if (handle_omp_array_sections (c))
remove = true;
else
Index: gcc/omp-low.c
===================================================================
--- gcc/omp-low.c (revision 226462)
+++ gcc/omp-low.c (working copy)
@@ -1172,14 +1172,12 @@ build_outer_var_ref (tree var, omp_conte
if (gimple_code (ctx->stmt) == GIMPLE_OMP_FOR
&& gimple_omp_for_kind (ctx->stmt) == GF_OMP_FOR_KIND_OACC_LOOP)
{
- for (ctx = ctx->outer; ctx && !maybe_lookup_decl (var, ctx);
- ctx = ctx->outer)
- ;
-
- if (ctx == NULL)
- gcc_unreachable ();
-
- x = lookup_decl (var, ctx);
+ do
+ {
+ ctx = ctx->outer;
+ x = maybe_lookup_decl (var, ctx);
+ }
+ while(!x);
}
else
x = lookup_decl (var, ctx->outer);
@@ -1848,10 +1846,6 @@ scan_sharing_clauses (tree clauses, omp_
/* FALLTHRU */
case OMP_CLAUSE_FIRSTPRIVATE:
- if (is_gimple_omp_oacc (ctx->stmt))
- /* Clause represented by a gang-local map under OpenACC. */
- gcc_unreachable ();
- /* FALLTHRU */
case OMP_CLAUSE_REDUCTION:
case OMP_CLAUSE_LINEAR:
decl = OMP_CLAUSE_DECL (c);
@@ -1879,10 +1873,20 @@ scan_sharing_clauses (tree clauses, omp_
else if (!global)
install_var_field (decl, by_ref, 3, ctx);
}
- /* The gimplifier always includes a OMP_CLAUSE_MAP with each parallel
- reduction variable. So don't install a local variable here. */
+
if (!is_oacc_parallel (ctx))
install_var_local (decl, ctx);
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
+ {
+ install_var_field (decl, (TREE_CODE (TREE_TYPE (decl))
+ != REFERENCE_TYPE), 3, ctx);
+ install_var_local (decl, ctx);
+ }
+ else
+ /* The gimplifier always includes a OMP_CLAUSE_MAP with
+ each parallel reduction variable. So don't install a
+ local variable here. */
+ gcc_assert (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION);
break;
case OMP_CLAUSE__LOOPTEMP_:
@@ -2063,12 +2067,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_REDUCTION:
case OMP_CLAUSE_LINEAR:
@@ -11712,7 +11710,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;
- gimple_seq tgt_body, olist, ilist, orlist, irlist, new_body;
+ gimple_seq tgt_body, olist, ilist, orlist, irlist, fplist, new_body;
location_t loc = gimple_location (stmt);
bool offloaded, data_region, has_reduction;
unsigned int map_cnt = 0;
@@ -11764,6 +11762,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))
@@ -11772,6 +11771,11 @@ lower_omp_target (gimple_stmt_iterator *
default:
break;
+ case OMP_CLAUSE_FIRSTPRIVATE:
+ if (is_oacc_parallel (ctx))
+ goto first_private;
+ break;
+
case OMP_CLAUSE_MAP:
#ifdef ENABLE_CHECKING
/* First check what we're prepared to handle in the following. */
@@ -11803,6 +11807,8 @@ lower_omp_target (gimple_stmt_iterator *
/* FALLTHRU */
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
+ first_private:
+
var = OMP_CLAUSE_DECL (c);
if (!DECL_P (var))
{
@@ -11829,11 +11835,26 @@ 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);
- if (DECL_P (new_var))
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
+ {
+ if (TREE_CODE (TREE_TYPE (new_var)) == REFERENCE_TYPE)
+ {
+ /* 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;
@@ -11856,6 +11877,7 @@ lower_omp_target (gimple_stmt_iterator *
}
}
map_cnt++;
+ break;
}
if (offloaded)
@@ -11945,6 +11967,10 @@ lower_omp_target (gimple_stmt_iterator *
default:
break;
+ case OMP_CLAUSE_FIRSTPRIVATE:
+ if (!is_oacc_parallel (ctx))
+ break;
+ /* FALLTHROUGH */
case OMP_CLAUSE_MAP:
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
@@ -12011,6 +12037,14 @@ 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)
+ {
+ if (TREE_CODE (TREE_TYPE (var)) != REFERENCE_TYPE)
+ 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);
@@ -12039,7 +12073,16 @@ lower_omp_target (gimple_stmt_iterator *
gimplify_assign (x, var, &ilist);
}
}
- tree s = OMP_CLAUSE_SIZE (c);
+ tree s = NULL_TREE;
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_FIRSTPRIVATE)
+ s = OMP_CLAUSE_SIZE (c);
+ else
+ {
+ s = TREE_TYPE (ovar);
+ if (TREE_CODE (s) == REFERENCE_TYPE)
+ s = TREE_TYPE (s);
+ s = TYPE_SIZE_UNIT (s);
+ }
if (s == NULL_TREE)
s = TYPE_SIZE_UNIT (TREE_TYPE (ovar));
s = fold_convert (size_type_node, s);
@@ -12054,6 +12097,9 @@ lower_omp_target (gimple_stmt_iterator *
case OMP_CLAUSE_MAP:
tkind = OMP_CLAUSE_MAP_KIND (c);
break;
+ case OMP_CLAUSE_FIRSTPRIVATE:
+ tkind = GOMP_MAP_TO;
+ break;
case OMP_CLAUSE_TO:
tkind = GOMP_MAP_TO;
break;
@@ -12118,6 +12164,7 @@ lower_omp_target (gimple_stmt_iterator *
gimple_build_assign (ctx->receiver_decl, t));
}
gimple_seq_add_seq (&new_body, ctx->ganglocal_init);
+ gimple_seq_add_seq (&new_body, fplist);
if (offloaded)
{
Index: gcc/fortran/trans-openmp.c
===================================================================
--- gcc/fortran/trans-openmp.c (revision 226462)
+++ gcc/fortran/trans-openmp.c (working copy)
@@ -2125,9 +2125,6 @@ gfc_trans_omp_clauses_1 (stmtblock_t *bl
case OMP_MAP_FROM:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FROM);
break;
- case OMP_MAP_GANGLOCAL:
- OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_TO_GANGLOCAL);
- break;
case OMP_MAP_TOFROM:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_TOFROM);
break;
@@ -2152,9 +2149,6 @@ gfc_trans_omp_clauses_1 (stmtblock_t *bl
case OMP_MAP_FORCE_DEVICEPTR:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_DEVICEPTR);
break;
- case OMP_MAP_FORCE_TO_GANGLOCAL:
- OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_TO_GANGLOCAL);
- break;
case OMP_MAP_DEVICE_RESIDENT:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_DEVICE_RESIDENT);
break;
Index: gcc/fortran/gfortran.h
===================================================================
--- gcc/fortran/gfortran.h (revision 226462)
+++ gcc/fortran/gfortran.h (working copy)
@@ -1138,7 +1138,6 @@ typedef enum
OMP_MAP_ALLOC,
OMP_MAP_TO,
OMP_MAP_FROM,
- OMP_MAP_GANGLOCAL,
OMP_MAP_TOFROM,
OMP_MAP_FORCE_ALLOC,
OMP_MAP_FORCE_DEALLOC,
@@ -1149,7 +1148,6 @@ typedef enum
OMP_MAP_FORCE_DEVICEPTR,
OMP_MAP_DEVICE_RESIDENT,
OMP_MAP_LINK,
- OMP_MAP_FORCE_TO_GANGLOCAL
}
gfc_omp_map_op;
Index: gcc/fortran/openmp.c
===================================================================
--- gcc/fortran/openmp.c (revision 226462)
+++ gcc/fortran/openmp.c (working copy)
@@ -586,22 +586,12 @@ gfc_match_omp_clauses (gfc_omp_clauses *
&c->lists[OMP_LIST_PRIVATE], true)
== MATCH_YES)
continue;
- if (mask & OMP_CLAUSE_FIRSTPRIVATE)
- {
- if (openacc)
- {
- if (gfc_match ("firstprivate ( ") == MATCH_YES
- && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_GANGLOCAL, false))
- continue;
- }
- else if (gfc_match_omp_variable_list ("firstprivate (",
+ if ((mask & OMP_CLAUSE_FIRSTPRIVATE)
+ && gfc_match_omp_variable_list ("firstprivate (",
&c->lists[OMP_LIST_FIRSTPRIVATE],
- true)
- == MATCH_YES)
- continue;
-
- }
+ true)
+ == MATCH_YES)
+ continue;
if ((mask & OMP_CLAUSE_LASTPRIVATE)
&& gfc_match_omp_variable_list ("lastprivate (",
&c->lists[OMP_LIST_LASTPRIVATE],
Index: gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 (revision 226462)
+++ gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 (working copy)
@@ -37,4 +37,3 @@ end program test
! { dg-final { scan-tree-dump-times "map\\(force_deviceptr:u\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "private\\(v\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "map\\(force_to_ganglocal:w" 1 "original" } }
Index: gcc/testsuite/gfortran.dg/goacc/list.f95
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/list.f95 (revision 226462)
+++ gcc/testsuite/gfortran.dg/goacc/list.f95 (working copy)
@@ -5,7 +5,7 @@ program test
implicit none
integer :: i, j, k, l, a(10)
- common /b/ j, k
+ common /b/ k
real, pointer :: p1 => NULL()
complex :: c, d(10)
@@ -64,8 +64,8 @@ program test
!$acc parallel firstprivate(10) ! { dg-error "Syntax error" }
- !$acc parallel firstprivate (/b/, /b/) ! { dg-error "Syntax error" }
- !$acc end parallel ! { dg-error "Unexpected" }
+ !$acc parallel firstprivate (/b/, /b/) ! { dg-error "present on multiple clauses" }
+ !$acc end parallel
!$acc parallel firstprivate (i, j, i) ! { dg-error "present on multiple clauses" }
!$acc end parallel
Index: gcc/testsuite/c-c++-common/goacc/firstprivate.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/firstprivate.c (revision 226462)
+++ gcc/testsuite/c-c++-common/goacc/firstprivate.c (working copy)
@@ -4,6 +4,6 @@ foo (void)
int a, b[100];
#pragma acc parallel firstprivate (a, b)
;
-#pragma acc parallel firstprivate (b[10:20]) /* { dg-error "subarrays are not permitted in firstprivate" } */
+#pragma acc parallel firstprivate (b[10:20]) /* { dg-error "expected" } */
;
}
Index: gcc/cp/semantics.c
===================================================================
--- gcc/cp/semantics.c (revision 226462)
+++ gcc/cp/semantics.c (working copy)
@@ -5838,10 +5838,6 @@ finish_omp_clauses (tree clauses, bool o
t = OMP_CLAUSE_DECL (c);
if (TREE_CODE (t) == TREE_LIST)
{
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_TO_GANGLOCAL))
- error_at (OMP_CLAUSE_LOCATION (c),
- "subarrays are not permitted in firstprivate");
if (handle_omp_array_sections (c))
remove = true;
else
Index: gcc/cp/parser.c
===================================================================
--- gcc/cp/parser.c (revision 226462)
+++ gcc/cp/parser.c (working copy)
@@ -28195,9 +28195,6 @@ cp_parser_oacc_data_clause (cp_parser *p
case PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT:
kind = GOMP_MAP_DEVICE_RESIDENT;
break;
- case PRAGMA_OACC_CLAUSE_FIRSTPRIVATE:
- kind = GOMP_MAP_FORCE_TO_GANGLOCAL;
- break;
case PRAGMA_OACC_CLAUSE_HOST:
kind = GOMP_MAP_FORCE_FROM;
break;
@@ -29753,7 +29750,8 @@ cp_parser_oacc_all_clauses (cp_parser *p
c_name = "deviceptr";
break;
case PRAGMA_OACC_CLAUSE_FIRSTPRIVATE:
- clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ clauses = cp_parser_omp_var_list
+ (parser, OMP_CLAUSE_FIRSTPRIVATE, clauses);
c_name = "firstprivate";
break;
case PRAGMA_OACC_CLAUSE_IF:
