This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
[gomp4] OpenACC first private
- From: Nathan Sidwell <nathan at acm dot org>
- To: GCC Patches <gcc-patches at gcc dot gnu dot org>
- Cc: james norris <James_Norris at mentor dot com>
- Date: Mon, 3 Aug 2015 10:30:49 -0400
- Subject: [gomp4] OpenACC first private
- Authentication-results: sourceware.org; auth=none
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 <nathan@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: