* [gomp4] OpenACC first private
@ 2015-08-03 14:30 Nathan Sidwell
2015-08-10 19:50 ` Thomas Schwinge
` (2 more replies)
0 siblings, 3 replies; 8+ messages in thread
From: Nathan Sidwell @ 2015-08-03 14:30 UTC (permalink / raw)
To: GCC Patches; +Cc: james norris
[-- Attachment #1: Type: text/plain, Size: 498 bytes --]
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
[-- Attachment #2: gomp4-fp.patch --]
[-- Type: text/x-patch, Size: 17707 bytes --]
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:
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [gomp4] OpenACC first private
2015-08-03 14:30 [gomp4] OpenACC first private Nathan Sidwell
@ 2015-08-10 19:50 ` Thomas Schwinge
2015-08-18 23:30 ` Thomas Schwinge
2015-08-12 18:31 ` [gomp4] declare directive James Norris
2015-10-29 8:31 ` [gomp4] OpenACC first private Thomas Schwinge
2 siblings, 1 reply; 8+ messages in thread
From: Thomas Schwinge @ 2015-08-10 19:50 UTC (permalink / raw)
To: Nathan Sidwell, GCC Patches; +Cc: james norris
[-- Attachment #1: Type: text/plain, Size: 8261 bytes --]
Hi!
On Mon, 3 Aug 2015 10:30:49 -0400, Nathan Sidwell <nathan@acm.org> wrote:
> 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.
I'm seeing the following regressions after this r226508 commit -- are
those the ones that Jim is working on resolving?
[-PASS:-]{+FAIL: gfortran.dg/goacc/modules.f95 -O (internal compiler error)+}
{+FAIL:+} gfortran.dg/goacc/modules.f95 -O (test for excess errors)
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-loop-1.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 (test for excess errors)
[-XFAIL:-]{+XPASS:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-loop-1.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 execution test
PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-4.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 (test for excess errors)
{+WARNING: program timed out.+}
XFAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-4.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 execution test
PASS: libgomp.oacc-c++/../libgomp.oacc-c-c++-common/parallel-loop-1.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 (test for excess errors)
[-XFAIL:-]{+XPASS:+} libgomp.oacc-c++/../libgomp.oacc-c-c++-common/parallel-loop-1.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 execution test
PASS: libgomp.oacc-c++/../libgomp.oacc-c-c++-common/reduction-4.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 (test for excess errors)
{+WARNING: program timed out.+}
XFAIL: libgomp.oacc-c++/../libgomp.oacc-c-c++-common/reduction-4.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 execution test
PASS: libgomp.oacc-fortran/declare-1.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O0 (test for excess errors)
[-PASS:-]{+FAIL:+} libgomp.oacc-fortran/declare-1.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O0 execution test
PASS: libgomp.oacc-fortran/declare-1.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O1 (test for excess errors)
[-PASS:-]{+FAIL:+} libgomp.oacc-fortran/declare-1.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O1 execution test
PASS: libgomp.oacc-fortran/declare-1.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O2 (test for excess errors)
[-PASS:-]{+FAIL:+} libgomp.oacc-fortran/declare-1.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O2 execution test
PASS: libgomp.oacc-fortran/declare-1.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -fomit-frame-pointer (test for excess errors)
[-PASS:-]{+FAIL:+} libgomp.oacc-fortran/declare-1.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -fomit-frame-pointer execution test
PASS: libgomp.oacc-fortran/declare-1.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -fomit-frame-pointer -funroll-loops (test for excess errors)
[-PASS:-]{+FAIL:+} libgomp.oacc-fortran/declare-1.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -fomit-frame-pointer -funroll-loops execution test
PASS: libgomp.oacc-fortran/declare-1.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -fomit-frame-pointer -funroll-all-loops -finline-functions (test for excess errors)
[-PASS:-]{+FAIL:+} libgomp.oacc-fortran/declare-1.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -fomit-frame-pointer -funroll-all-loops -finline-functions execution test
PASS: libgomp.oacc-fortran/declare-1.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -g (test for excess errors)
[-PASS:-]{+FAIL:+} libgomp.oacc-fortran/declare-1.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -g execution test
PASS: libgomp.oacc-fortran/declare-1.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -Os (test for excess errors)
[-PASS:-]{+FAIL:+} libgomp.oacc-fortran/declare-1.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -Os execution test
PASS: libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O0 (test for excess errors)
[-PASS:-]{+FAIL:+} libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O0 execution test
PASS: libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O1 (test for excess errors)
[-PASS:-]{+FAIL:+} libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O1 execution test
PASS: libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O2 (test for excess errors)
PASS: libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O2 execution test
PASS: libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -fomit-frame-pointer (test for excess errors)
PASS: libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -fomit-frame-pointer execution test
PASS: libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -fomit-frame-pointer -funroll-loops (test for excess errors)
[-PASS:-]{+FAIL:+} libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -fomit-frame-pointer -funroll-loops execution test
PASS: libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -fomit-frame-pointer -funroll-all-loops -finline-functions (test for excess errors)
PASS: libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -fomit-frame-pointer -funroll-all-loops -finline-functions execution test
PASS: libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -g (test for excess errors)
> --- 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;
In file included from [...]/source-gcc/gcc/fortran/arith.c:30:0:
[...]/source-gcc/gcc/fortran/gfortran.h:1150:15: warning: comma at end of enumerator list [-Wpedantic]
OMP_MAP_LINK,
^
Committed to gomp-4_0-branch in r226768:
commit 0ab0693710e7f0f88b8966233c84295f4f6179fc
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Mon Aug 10 19:48:48 2015 +0000
Address -Wpedantic diagnostic
Fixup for r226508.
gcc/fortran/
* gfortran.h (gfc_omp_map_op): Remove comma at end of enumerator
list.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@226768 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/fortran/ChangeLog.gomp | 5 +++++
gcc/fortran/gfortran.h | 2 +-
2 files changed, 6 insertions(+), 1 deletion(-)
diff --git gcc/fortran/ChangeLog.gomp gcc/fortran/ChangeLog.gomp
index 21aa06e..fd7204c 100644
--- gcc/fortran/ChangeLog.gomp
+++ gcc/fortran/ChangeLog.gomp
@@ -1,3 +1,8 @@
+2015-08-10 Thomas Schwinge <thomas@codesourcery.com>
+
+ * gfortran.h (gfc_omp_map_op): Remove comma at end of enumerator
+ list.
+
2015-08-03 Nathan Sidwell <nathan@codesourcery.com>
* trans-openmp.c (gfc_trans_omp_clauses_1): Remove GANGLOCAL
diff --git gcc/fortran/gfortran.h gcc/fortran/gfortran.h
index 1e87d6c..128d65e 100644
--- gcc/fortran/gfortran.h
+++ gcc/fortran/gfortran.h
@@ -1147,7 +1147,7 @@ typedef enum
OMP_MAP_FORCE_PRESENT,
OMP_MAP_FORCE_DEVICEPTR,
OMP_MAP_DEVICE_RESIDENT,
- OMP_MAP_LINK,
+ OMP_MAP_LINK
}
gfc_omp_map_op;
Grüße,
Thomas
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]
^ permalink raw reply [flat|nested] 8+ messages in thread
* [gomp4] declare directive
@ 2015-08-12 18:31 ` James Norris
2015-09-01 16:34 ` Tom de Vries
2015-09-02 8:25 ` Tom de Vries
0 siblings, 2 replies; 8+ messages in thread
From: James Norris @ 2015-08-12 18:31 UTC (permalink / raw)
To: gcc-patches; +Cc: Jakub Jelinek
[-- Attachment #1: Type: text/plain, Size: 3347 bytes --]
Hi,
The attached patch is a revision of the functionality required to
handle the declare directive. The motivation for the change was that
the original code was deemed to fragile and difficult to maintain.
The new code is now smaller, in terms of line count, and hopefully,
easier to understand. Given the significant amount of changes, I've
included a commentary on entire functions, rather than comments
about the specific changes. Hopefully, this will help in the
review process.
gimple.c : gimplify_oacc_declare ()
The first part of this function deals with handling
the entry clauses. After scanning the clauses, we
spin through the list and while adding the appropriate
ones to the context variable list. Also if any
'oacc declare' attributes are around, toss them.
The next section deals with the return clauses.
Again scan them, but after scanning them toss the
context that was created as a result of the scan.
We don't need the context as the context that was
created with the entry clauses will suffice. After
creation of the gimple statement, save it away so
that it can be used at exit time.
gimple.c : gimplify_body ()
This is the place where the 'exit' clauses are
used and inserted via a gimple statement. This
is also the place where the omp context is
'unwound'.
fortran/trans-decl.c: finish_oacc_declare ()
This function is called to handle the declare
directives. The information that was contained
in the directives during parsing of the variable
section is hanging off the namespace (ns).
The primary function here is to create the two
sets of derived clauses and associate them with
a newly created statement, i.e, gfc_code. This
new statement is then inserted at the beginning
of the statement chain associated with 'ns'.
c/c-parser.c c_parser_oacc_declare ()
The initial section of the function is doing
syntax checking. If no errors are found and
the variable is 'global', then an attribute
is added to the variable. This is followed by
the creation of the derived clauses.
In the final section, when a global is at hand,
a constructor is created. The constructor will
pass on information to the runtime. On the other
hand if the variable is local, then a statement
will be inserted and created that will contain
both sets of derived clauses.
cp/parser.c: cp_parser_oacc_declare ()
The initial section of the function is doing
syntax checking. If no errors are found and
the variable is 'global', then an attribute
is added to the variable. This is followed by
the creation of the derived clauses.
In the final section, when a global is at hand,
a constructor is created. The constructor will
pass on information to the runtime. On the other
hand if the variable is local, then a statement
will be inserted and created that will contain
both sets of derived clauses.
I'll commit this to gomp-4_0-branch sometime during the week of 17 August.
Thanks!
Jim
[-- Attachment #2: declare.patch --]
[-- Type: text/x-patch, Size: 42416 bytes --]
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index b452235..3f3e8c0 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -1452,319 +1452,6 @@ c_parser_external_declaration (c_parser *parser)
}
}
-static tree
-check_oacc_vars_1 (tree *tp, int *, void *l)
-{
- if (TREE_CODE (*tp) == VAR_DECL && TREE_PUBLIC (*tp))
- {
- location_t loc = DECL_SOURCE_LOCATION (*tp);
- tree attrs;
- attrs = lookup_attribute ("oacc declare", DECL_ATTRIBUTES (*tp));
- if (attrs)
- {
- tree t;
-
- for (t = TREE_VALUE (attrs); t; t = TREE_CHAIN (t))
- {
- loc = EXPR_LOCATION ((tree) l);
-
- if (OMP_CLAUSE_MAP_KIND (TREE_VALUE (t)) == GOMP_MAP_LINK)
- {
- error_at (loc, "%<link%> clause cannot be used with %qE",
- *tp);
- break;
- }
- }
- }
- else
- error_at (loc, "no %<#pragma acc declare%> for %qE", *tp);
- }
- return NULL_TREE;
-}
-
-static tree
-check_oacc_vars (tree *tp, int *, void *)
-{
- if (TREE_CODE (*tp) == STATEMENT_LIST)
- {
- tree_stmt_iterator i;
-
- for (i = tsi_start (*tp); !tsi_end_p (i); tsi_next (&i))
- {
- tree t = tsi_stmt (i);
- walk_tree_without_duplicates (&t, check_oacc_vars_1, t);
- }
- }
-
- return NULL_TREE;
-}
-
-static struct oacc_return
-{
- tree_stmt_iterator iter;
- tree stmt;
- int op;
- struct oacc_return *next;
-} *oacc_returns;
-
-static tree
-find_oacc_return (tree *tp, int *, void *)
-{
- if (TREE_CODE (*tp) == STATEMENT_LIST)
- {
- tree_stmt_iterator i;
-
- for (i = tsi_start (*tp); !tsi_end_p (i); tsi_next (&i))
- {
- tree t;
- struct oacc_return *r;
-
- t = tsi_stmt (i);
-
- if (TREE_CODE (t) == RETURN_EXPR)
- {
- r = XNEW (struct oacc_return);
- r->iter = i;
- r->stmt = NULL_TREE;
- r->op = 1;
- r->next = NULL;
-
- if (oacc_returns)
- r->next = oacc_returns;
-
- oacc_returns = r;
- }
- else if (TREE_CODE (t) == COND_EXPR)
- {
- bool op1, op2;
- tree op;
-
- op1 = op2 = false;
-
- op = TREE_OPERAND (t, 1);
- op1 = (op && TREE_CODE (op) == RETURN_EXPR);
-
- op = TREE_OPERAND (t, 2);
- op2 = (op && TREE_CODE (op) == RETURN_EXPR);
-
- if (op1 || op2)
- {
- r = XNEW (struct oacc_return);
- r->stmt = t;
- r->op = op1 ? 1 : 2;
- r->next = NULL;
-
- if (oacc_returns)
- r->next = oacc_returns;
-
- oacc_returns = r;
- }
- }
- }
- }
-
- return NULL_TREE;
-}
-
-static void
-finish_oacc_declare (tree fnbody, tree decls)
-{
- tree t, stmt, body, c, ret_clauses, clauses;
- location_t loc;
- tree_stmt_iterator i;
- tree fndecl = current_function_decl;
-
- if (lookup_attribute ("oacc function", DECL_ATTRIBUTES (fndecl)))
- {
- if (lookup_attribute ("oacc declare", DECL_ATTRIBUTES (fndecl)))
- {
- location_t loc = DECL_SOURCE_LOCATION (fndecl);
- error_at (loc, "%<#pragma acc declare%> not allowed in %qE", fndecl);
- }
-
- walk_tree_without_duplicates (&fnbody, check_oacc_vars, NULL);
- return;
- }
-
- if (!decls)
- return;
-
- body = BIND_EXPR_BODY (fnbody);
-
- if (TREE_CODE (body) != STATEMENT_LIST)
- {
- tree list;
-
- list = alloc_stmt_list ();
- append_to_statement_list (body, &list);
- BIND_EXPR_BODY (fnbody) = list;
- body = list;
- }
-
- walk_tree_without_duplicates (&body, find_oacc_return, NULL);
-
- clauses = NULL_TREE;
-
- for (t = decls; t; t = TREE_CHAIN (t))
- {
- c = TREE_VALUE (TREE_VALUE (t));
-
- if (clauses)
- OMP_CLAUSE_CHAIN (c) = clauses;
- else
- loc = OMP_CLAUSE_LOCATION (c);
-
- clauses = c;
- }
-
- ret_clauses = NULL_TREE;
-
- for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
- {
- bool ret = false;
- HOST_WIDE_INT kind, new_op;
-
- kind = OMP_CLAUSE_MAP_KIND (c);
-
- switch (kind)
- {
- case GOMP_MAP_ALLOC:
- case GOMP_MAP_FORCE_ALLOC:
- case GOMP_MAP_FORCE_TO:
- new_op = GOMP_MAP_FORCE_DEALLOC;
- ret = true;
- break;
-
- case GOMP_MAP_FORCE_FROM:
- OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_ALLOC);
- new_op = GOMP_MAP_FORCE_FROM;
- ret = true;
- break;
-
- case GOMP_MAP_FORCE_TOFROM:
- OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_TO);
- new_op = GOMP_MAP_FORCE_FROM;
- ret = true;
- break;
-
- case GOMP_MAP_FROM:
- OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_ALLOC);
- new_op = GOMP_MAP_FROM;
- ret = true;
- break;
-
- case GOMP_MAP_TOFROM:
- OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TO);
- new_op = GOMP_MAP_FROM;
- ret = true;
- break;
-
- case GOMP_MAP_DEVICE_RESIDENT:
- case GOMP_MAP_FORCE_DEVICEPTR:
- case GOMP_MAP_FORCE_PRESENT:
- case GOMP_MAP_LINK:
- case GOMP_MAP_POINTER:
- case GOMP_MAP_TO:
- break;
-
- default:
- gcc_unreachable ();
- break;
- }
-
- if (ret)
- {
- t = copy_node (c);
-
- OMP_CLAUSE_SET_MAP_KIND (t, new_op);
-
- if (ret_clauses)
- OMP_CLAUSE_CHAIN (t) = ret_clauses;
-
- ret_clauses = t;
- }
- }
-
- if (clauses)
- {
- bool found = false;
-
- stmt = make_node (OACC_DECLARE);
- TREE_TYPE (stmt) = void_type_node;
- OACC_DECLARE_CLAUSES (stmt) = clauses;
- SET_EXPR_LOCATION (stmt, loc);
-
- c = OMP_CLAUSE_DECL (TREE_VALUE (TREE_VALUE (decls)));
-
- for (i = tsi_start (body); !tsi_end_p (i); tsi_next (&i))
- {
- tree it;
-
- it = tsi_stmt (i);
-
- if ((TREE_CODE (it) == DECL_EXPR) && (DECL_EXPR_DECL (it) == c))
- {
- tsi_link_after (&i, stmt, TSI_CONTINUE_LINKING);
- found = true;
- break;
- }
- }
-
- if (!found)
- {
- i = tsi_start (body);
- tsi_link_before (&i, stmt, TSI_CONTINUE_LINKING);
- }
- }
-
- while (oacc_returns)
- {
- struct oacc_return *r;
-
- stmt = make_node (OACC_DECLARE);
- TREE_TYPE (stmt) = void_type_node;
- OACC_DECLARE_CLAUSES (stmt) = ret_clauses;
- SET_EXPR_LOCATION (stmt, loc);
-
- r = oacc_returns;
- if (r->stmt)
- {
- tree l;
-
- l = alloc_stmt_list ();
- append_to_statement_list (stmt, &l);
- stmt = TREE_OPERAND (r->stmt, r->op);
- append_to_statement_list (stmt, &l);
- TREE_OPERAND (r->stmt, r->op) = l;
- }
- else
- tsi_link_before (&r->iter, stmt, TSI_CONTINUE_LINKING);
-
- oacc_returns = r->next;
- free (r);
- }
-
- for (i = tsi_start (body); !tsi_end_p (i); tsi_next (&i))
- {
- if (tsi_end_p (i))
- break;
- }
-
- if (ret_clauses)
- {
- stmt = make_node (OACC_DECLARE);
- TREE_TYPE (stmt) = void_type_node;
- OACC_DECLARE_CLAUSES (stmt) = ret_clauses;
- SET_EXPR_LOCATION (stmt, loc);
-
- tsi_link_before (&i, stmt, TSI_CONTINUE_LINKING);
- }
-
- DECL_ATTRIBUTES (fndecl)
- = remove_attribute ("oacc declare", DECL_ATTRIBUTES (fndecl));
-}
-
-
static void c_finish_omp_declare_simd (c_parser *, tree, tree, vec<c_token>);
static void c_finish_oacc_routine (c_parser *, tree, tree, bool);
@@ -2312,9 +1999,6 @@ c_parser_declaration_or_fndef (c_parser *parser, bool fndef_ok,
fnbody = c_parser_compound_statement (parser);
if (flag_cilkplus && contains_array_notation_expr (fnbody))
fnbody = expand_array_notation_exprs (fnbody);
- tree decls = lookup_attribute ("oacc declare",
- DECL_ATTRIBUTES (current_function_decl));
- finish_oacc_declare (fnbody, decls);
if (nested)
{
tree decl = current_function_decl;
@@ -12770,7 +12454,8 @@ static void
c_parser_oacc_declare (c_parser *parser)
{
location_t pragma_loc = c_parser_peek_token (parser)->location;
- tree clauses;
+ tree c, clauses, ret_clauses, stmt, t;
+
bool error = false;
c_parser_consume_pragma (parser);
@@ -12783,7 +12468,8 @@ c_parser_oacc_declare (c_parser *parser)
"no valid clauses specified in %<#pragma acc declare%>");
return;
}
- for (tree t = clauses; t; t = OMP_CLAUSE_CHAIN (t))
+
+ for (t = clauses; t; t = OMP_CLAUSE_CHAIN (t))
{
location_t loc = OMP_CLAUSE_LOCATION (t);
tree decl = OMP_CLAUSE_DECL (t);
@@ -12849,60 +12535,27 @@ c_parser_oacc_declare (c_parser *parser)
break;
}
- /* Store the clause in an attribute on the variable, at file
- scope, or the function, at block scope. */
- tree decl_for_attr;
- if (global_bindings_p ())
+ tree decl_for_attr = decl;
+ tree prev_attr = lookup_attribute ("oacc declare",
+ DECL_ATTRIBUTES (decl));
+ if (prev_attr)
{
- decl_for_attr = decl;
- tree prev_attr = lookup_attribute ("oacc declare",
- DECL_ATTRIBUTES (decl));
- if (prev_attr)
- {
- tree p = TREE_VALUE (prev_attr);
- tree cl = TREE_VALUE (p);
+ tree p = TREE_VALUE (prev_attr);
+ tree cl = TREE_VALUE (p);
- if (!devres
- && OMP_CLAUSE_MAP_KIND (cl) != GOMP_MAP_DEVICE_RESIDENT)
- {
- error_at (loc,
- "variable %qD used more than once with "
- "%<#pragma acc declare%>", decl);
- inform (OMP_CLAUSE_LOCATION (cl),
- "previous directive was here");
- error = true;
- continue;
- }
- }
- }
- else
- {
- decl_for_attr = current_function_decl;
- tree prev_attr = lookup_attribute ("oacc declare",
- DECL_ATTRIBUTES (decl_for_attr));
- for (;
- prev_attr;
- prev_attr = lookup_attribute ("oacc declare",
- TREE_CHAIN (prev_attr)))
+ if (!devres && OMP_CLAUSE_MAP_KIND (cl) != GOMP_MAP_DEVICE_RESIDENT)
{
- tree p = TREE_VALUE (prev_attr);
- tree cl = TREE_VALUE (p);
- if (OMP_CLAUSE_DECL (cl) == decl)
- {
- error_at (loc,
- "variable %qD used more than once with "
- "%<#pragma acc declare%>", decl);
- inform (OMP_CLAUSE_LOCATION (cl),
- "previous directive was here");
- error = true;
- break;
- }
+ error_at (loc, "variable %qD used more than once with "
+ "%<#pragma acc declare%>", decl);
+ inform (OMP_CLAUSE_LOCATION (cl), "previous directive was here");
+ error = true;
+ continue;
}
}
- if (!error)
+ if (!error && global_bindings_p ())
{
- tree attr = tree_cons (NULL_TREE, t, NULL_TREE);
+ tree attr = tree_cons (NULL_TREE, clauses, NULL_TREE);
tree attrs = tree_cons (get_identifier ("oacc declare"),
attr, NULL_TREE);
decl_attributes (&decl_for_attr, attrs, 0);
@@ -12912,6 +12565,74 @@ c_parser_oacc_declare (c_parser *parser)
if (error)
return;
+ ret_clauses = NULL_TREE;
+
+ for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+ {
+ bool ret = false;
+ HOST_WIDE_INT kind, new_op;
+
+ kind = OMP_CLAUSE_MAP_KIND (c);
+
+ switch (kind)
+ {
+ case GOMP_MAP_ALLOC:
+ case GOMP_MAP_FORCE_ALLOC:
+ case GOMP_MAP_FORCE_TO:
+ new_op = GOMP_MAP_FORCE_DEALLOC;
+ ret = true;
+ break;
+
+ case GOMP_MAP_FORCE_FROM:
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_ALLOC);
+ new_op = GOMP_MAP_FORCE_FROM;
+ ret = true;
+ break;
+
+ case GOMP_MAP_FORCE_TOFROM:
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_TO);
+ new_op = GOMP_MAP_FORCE_FROM;
+ ret = true;
+ break;
+
+ case GOMP_MAP_FROM:
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_ALLOC);
+ new_op = GOMP_MAP_FROM;
+ ret = true;
+ break;
+
+ case GOMP_MAP_TOFROM:
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TO);
+ new_op = GOMP_MAP_FROM;
+ ret = true;
+ break;
+
+ case GOMP_MAP_DEVICE_RESIDENT:
+ case GOMP_MAP_FORCE_DEVICEPTR:
+ case GOMP_MAP_FORCE_PRESENT:
+ case GOMP_MAP_LINK:
+ case GOMP_MAP_POINTER:
+ case GOMP_MAP_TO:
+ break;
+
+ default:
+ gcc_unreachable ();
+ break;
+ }
+
+ if (ret)
+ {
+ t = copy_node (c);
+
+ OMP_CLAUSE_SET_MAP_KIND (t, new_op);
+
+ if (ret_clauses)
+ OMP_CLAUSE_CHAIN (t) = ret_clauses;
+
+ ret_clauses = t;
+ }
+ }
+
if (global_bindings_p ())
{
char buf[128];
@@ -12971,6 +12692,16 @@ c_parser_oacc_declare (c_parser *parser)
finish_function ();
}
+ else
+ {
+ stmt = make_node (OACC_DECLARE);
+ TREE_TYPE (stmt) = void_type_node;
+ OACC_DECLARE_CLAUSES (stmt) = clauses;
+ OACC_DECLARE_RETURN_CLAUSES (stmt) = ret_clauses;
+ SET_EXPR_LOCATION (stmt, pragma_loc);
+
+ add_stmt (stmt);
+ }
}
/* OpenACC 2.0:
diff --git a/gcc/cp/decl.c b/gcc/cp/decl.c
index 1d80ef2..069db46 100644
--- a/gcc/cp/decl.c
+++ b/gcc/cp/decl.c
@@ -14114,332 +14114,6 @@ maybe_save_function_definition (tree fun)
register_constexpr_fundef (fun, DECL_SAVED_TREE (fun));
}
-static tree
-check_oacc_vars_1 (tree *tp, int *, void *l)
-{
- if (TREE_CODE (*tp) == VAR_DECL && TREE_PUBLIC (*tp))
- {
- location_t loc = DECL_SOURCE_LOCATION (*tp);
- tree attrs;
- attrs = lookup_attribute ("oacc declare", DECL_ATTRIBUTES (*tp));
- if (attrs)
- {
- tree t;
-
- for (t = TREE_VALUE (attrs); t; t = TREE_CHAIN (t))
- {
- loc = EXPR_LOCATION ((tree) l);
-
- if (OMP_CLAUSE_MAP_KIND (TREE_VALUE (t)) == GOMP_MAP_LINK)
- {
- error_at (loc, "%<link%> clause cannot be used with %qE",
- *tp);
- break;
- }
- }
- }
- else
- error_at (loc, "no %<#pragma acc declare%> for %qE", *tp);
- }
- return NULL_TREE;
-}
-
-static tree
-check_oacc_vars (tree *tp, int *, void *)
-{
- if (TREE_CODE (*tp) == STATEMENT_LIST)
- {
- tree_stmt_iterator i;
-
- for (i = tsi_start (*tp); !tsi_end_p (i); tsi_next (&i))
- {
- tree t = tsi_stmt (i);
- walk_tree_without_duplicates (&t, check_oacc_vars_1, t);
- }
- }
-
- return NULL_TREE;
-}
-
-static struct oacc_return
-{
- tree_stmt_iterator iter;
- tree stmt;
- int op;
- struct oacc_return *next;
-} *oacc_returns;
-
-static tree
-find_oacc_return (tree *tp, int *, void *)
-{
- if (TREE_CODE (*tp) == STATEMENT_LIST)
- {
- tree_stmt_iterator i;
-
- for (i = tsi_start (*tp); !tsi_end_p (i); tsi_next (&i))
- {
- tree t;
- struct oacc_return *r;
-
- t = tsi_stmt (i);
-
- if (TREE_CODE (t) == RETURN_EXPR)
- {
- r = XNEW (struct oacc_return);
- r->iter = i;
- r->stmt = NULL_TREE;
- r->op = 1;
- r->next = NULL;
-
- if (oacc_returns)
- r->next = oacc_returns;
-
- oacc_returns = r;
- }
- else if (TREE_CODE (t) == IF_STMT)
- {
- bool op1, op2;
- tree op;
-
- op1 = op2 = false;
-
- op = TREE_OPERAND (t, 1);
- op1 = (op && TREE_CODE (op) == RETURN_EXPR);
-
- op = TREE_OPERAND (t, 2);
- op2 = (op && TREE_CODE (op) == RETURN_EXPR);
-
- if (op1 || op2)
- {
- r = XNEW (struct oacc_return);
- r->stmt = t;
- r->op = op1 ? 1 : 2;
- r->next = NULL;
-
- if (oacc_returns)
- r->next = oacc_returns;
-
- oacc_returns = r;
- }
- }
- }
- }
-
- return NULL_TREE;
-}
-
-static void
-finish_oacc_declare (tree fndecl)
-{
- tree t, stmt, list, c, ret_clauses, clauses, decls;
- location_t loc;
- tree_stmt_iterator i;
-
- if (DECL_USE_TEMPLATE (fndecl))
- return;
-
- list = cur_stmt_list;
-
- decls = lookup_attribute ("oacc declare", DECL_ATTRIBUTES (fndecl));
-
- if (lookup_attribute ("oacc function", DECL_ATTRIBUTES (fndecl)))
- {
- if (decls)
- {
- location_t loc = DECL_SOURCE_LOCATION (fndecl);
- error_at (loc, "%<#pragma acc declare%> not allowed in %qE", fndecl);
- }
-
- walk_tree_without_duplicates (&list, check_oacc_vars, NULL);
- return;
- }
-
- if (!decls)
- return;
-
- walk_tree_without_duplicates (&list, find_oacc_return, NULL);
-
- clauses = NULL_TREE;
-
- for (t = decls; t; t = TREE_CHAIN (t))
- {
- c = TREE_VALUE (TREE_VALUE (t));
-
- if (clauses)
- OMP_CLAUSE_CHAIN (c) = clauses;
- else
- loc = OMP_CLAUSE_LOCATION (c);
-
- clauses = c;
- }
-
- ret_clauses = NULL_TREE;
-
- for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
- {
- bool ret = false;
- HOST_WIDE_INT kind, new_op;
-
- kind = OMP_CLAUSE_MAP_KIND (c);
-
- switch (kind)
- {
- case GOMP_MAP_ALLOC:
- case GOMP_MAP_FORCE_ALLOC:
- case GOMP_MAP_FORCE_TO:
- new_op = GOMP_MAP_FORCE_DEALLOC;
- ret = true;
- break;
-
- case GOMP_MAP_FORCE_FROM:
- OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_ALLOC);
- new_op = GOMP_MAP_FORCE_FROM;
- ret = true;
- break;
-
- case GOMP_MAP_FORCE_TOFROM:
- OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_TO);
- new_op = GOMP_MAP_FORCE_FROM;
- ret = true;
- break;
-
- case GOMP_MAP_FROM:
- OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_ALLOC);
- new_op = GOMP_MAP_FROM;
- ret = true;
- break;
-
- case GOMP_MAP_TOFROM:
- OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TO);
- new_op = GOMP_MAP_FROM;
- ret = true;
- break;
-
- case GOMP_MAP_DEVICE_RESIDENT:
- case GOMP_MAP_FORCE_DEVICEPTR:
- case GOMP_MAP_FORCE_PRESENT:
- case GOMP_MAP_POINTER:
- case GOMP_MAP_TO:
- break;
-
- case GOMP_MAP_LINK:
- continue;
-
- default:
- gcc_unreachable ();
- break;
- }
-
- if (ret)
- {
- t = copy_node (c);
-
- OMP_CLAUSE_SET_MAP_KIND (t, new_op);
-
- if (ret_clauses)
- OMP_CLAUSE_CHAIN (t) = ret_clauses;
-
- ret_clauses = t;
- }
- }
-
- i = tsi_start (list);
- if (!tsi_end_p (i))
- {
- t = tsi_stmt (i);
- if (TREE_CODE (t) == BIND_EXPR)
- {
- list = BIND_EXPR_BODY (t);
- if (TREE_CODE (list) != STATEMENT_LIST)
- {
- stmt = list;
- list = alloc_stmt_list ();
- BIND_EXPR_BODY (t) = list;
- i = tsi_start (list);
- tsi_link_after (&i, stmt, TSI_CONTINUE_LINKING);
- }
- }
- }
-
- if (clauses)
- {
- bool found = false;
-
- stmt = make_node (OACC_DECLARE);
- TREE_TYPE (stmt) = void_type_node;
- OMP_STANDALONE_CLAUSES (stmt) = clauses;
- SET_EXPR_LOCATION (stmt, loc);
-
- c = OMP_CLAUSE_DECL (TREE_VALUE (TREE_VALUE (decls)));
-
- for (i = tsi_start (list); !tsi_end_p (i); tsi_next (&i))
- {
- tree it;
-
- it = tsi_stmt (i);
-
- if ((TREE_CODE (it) == DECL_EXPR) && (DECL_EXPR_DECL (it) == c))
- {
- tsi_link_after (&i, stmt, TSI_CONTINUE_LINKING);
- found = true;
- break;
- }
- }
-
- if (!found)
- {
- i = tsi_start (list);
- tsi_link_before (&i, stmt, TSI_CONTINUE_LINKING);
- }
- }
-
- while (oacc_returns)
- {
- struct oacc_return *r;
-
- stmt = make_node (OACC_DECLARE);
- TREE_TYPE (stmt) = void_type_node;
- OMP_STANDALONE_CLAUSES (stmt) = ret_clauses;
- SET_EXPR_LOCATION (stmt, loc);
-
- r = oacc_returns;
- if (r->stmt)
- {
- tree l;
-
- l = alloc_stmt_list ();
- append_to_statement_list (stmt, &l);
- stmt = TREE_OPERAND (r->stmt, r->op);
- append_to_statement_list (stmt, &l);
- TREE_OPERAND (r->stmt, r->op) = l;
- }
- else
- tsi_link_before (&r->iter, stmt, TSI_CONTINUE_LINKING);
-
- oacc_returns = r->next;
- free (r);
- }
-
- if (ret_clauses)
- {
- for (i = tsi_start (list); !tsi_end_p (i); tsi_next (&i))
- {
- if (tsi_end_p (i))
- break;
- }
-
- stmt = make_node (OACC_DECLARE);
- TREE_TYPE (stmt) = void_type_node;
- OMP_STANDALONE_CLAUSES (stmt) = ret_clauses;
- SET_EXPR_LOCATION (stmt, loc);
-
- tsi_link_before (&i, stmt, TSI_CONTINUE_LINKING);
- }
-
- DECL_ATTRIBUTES (fndecl)
- = remove_attribute ("oacc declare", DECL_ATTRIBUTES (fndecl));
-}
-
/* Finish up a function declaration and compile that function
all the way to assembler language output. The free the storage
for the function definition.
@@ -14468,8 +14142,6 @@ finish_function (int flags)
gcc_assert (!defer_mark_used_calls);
defer_mark_used_calls = true;
- finish_oacc_declare (fndecl);
-
record_key_method_defined (fndecl);
fntype = TREE_TYPE (fndecl);
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 85236bf..14e7f8e 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -32179,12 +32179,14 @@ static int oacc_dcl_idx = 0;
static tree
cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
{
- tree clauses;
+ tree c, clauses, ret_clauses, stmt, t;
bool error = false;
+
clauses = cp_parser_oacc_all_clauses (parser, OACC_DECLARE_CLAUSE_MASK,
"#pragma acc declare", pragma_tok);
+
if (find_omp_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE)
{
error_at (pragma_tok->location,
@@ -32258,58 +32260,26 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
break;
}
- /* Store the clause in an attribute on the variable, at file
- scope, or the function, at block scope. */
- tree decl_for_attr;
- if (global_bindings_p ())
- {
- decl_for_attr = decl;
- tree prev_attr = lookup_attribute ("oacc declare",
+ tree decl_for_attr = decl;
+ tree prev_attr = lookup_attribute ("oacc declare",
DECL_ATTRIBUTES (decl));
- if (prev_attr)
- {
- tree p = TREE_VALUE (prev_attr);
- tree cl = TREE_VALUE (p);
-
- if (!devres
- && OMP_CLAUSE_MAP_KIND (cl) != GOMP_MAP_DEVICE_RESIDENT)
- {
- error_at (loc,
- "variable %qD used more than once with "
- "%<#pragma acc declare%>", decl);
- inform (OMP_CLAUSE_LOCATION (TREE_VALUE (p)),
- "previous directive was here");
- error = true;
- continue;
- }
- }
- }
- else
+ if (prev_attr)
{
- decl_for_attr = current_function_decl;
- tree prev_attr = lookup_attribute ("oacc declare",
- DECL_ATTRIBUTES (decl_for_attr));
- for (;
- prev_attr;
- prev_attr = lookup_attribute ("oacc declare",
- TREE_CHAIN (prev_attr)))
+ tree p = TREE_VALUE (prev_attr);
+ tree cl = TREE_VALUE (p);
+
+ if (!devres && OMP_CLAUSE_MAP_KIND (cl) != GOMP_MAP_DEVICE_RESIDENT)
{
- tree p = TREE_VALUE (prev_attr);
- tree cl = TREE_VALUE (p);
- if (OMP_CLAUSE_DECL (cl) == decl)
- {
- error_at (loc,
- "variable %qD used more than once with "
- "%<#pragma acc declare%>", decl);
- inform (OMP_CLAUSE_LOCATION (cl),
- "previous directive was here");
- error = true;
- break;
- }
+ error_at (loc, "variable %qD used more than once with "
+ "%<#pragma acc declare%>", decl);
+ inform (OMP_CLAUSE_LOCATION (TREE_VALUE (p)),
+ "previous directive was here");
+ error = true;
+ continue;
}
}
- if (!error)
+ if (!error && global_bindings_p ())
{
tree attr = tree_cons (NULL_TREE, t, NULL_TREE);
tree attrs = tree_cons (get_identifier ("oacc declare"),
@@ -32321,6 +32291,76 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
if (error)
return NULL_TREE;
+ ret_clauses = NULL_TREE;
+
+ for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+ {
+ bool ret = false;
+ HOST_WIDE_INT kind, new_op;
+
+ kind = OMP_CLAUSE_MAP_KIND (c);
+
+ switch (kind)
+ {
+ case GOMP_MAP_ALLOC:
+ case GOMP_MAP_FORCE_ALLOC:
+ case GOMP_MAP_FORCE_TO:
+ new_op = GOMP_MAP_FORCE_DEALLOC;
+ ret = true;
+ break;
+
+ case GOMP_MAP_FORCE_FROM:
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_ALLOC);
+ new_op = GOMP_MAP_FORCE_FROM;
+ ret = true;
+ break;
+
+ case GOMP_MAP_FORCE_TOFROM:
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_TO);
+ new_op = GOMP_MAP_FORCE_FROM;
+ ret = true;
+ break;
+
+ case GOMP_MAP_FROM:
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_ALLOC);
+ new_op = GOMP_MAP_FROM;
+ ret = true;
+ break;
+
+ case GOMP_MAP_TOFROM:
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TO);
+ new_op = GOMP_MAP_FROM;
+ ret = true;
+ break;
+
+ case GOMP_MAP_DEVICE_RESIDENT:
+ case GOMP_MAP_FORCE_DEVICEPTR:
+ case GOMP_MAP_FORCE_PRESENT:
+ case GOMP_MAP_POINTER:
+ case GOMP_MAP_TO:
+ break;
+
+ case GOMP_MAP_LINK:
+ continue;
+
+ default:
+ gcc_unreachable ();
+ break;
+ }
+
+ if (ret)
+ {
+ t = copy_node (c);
+
+ OMP_CLAUSE_SET_MAP_KIND (t, new_op);
+
+ if (ret_clauses)
+ OMP_CLAUSE_CHAIN (t) = ret_clauses;
+
+ ret_clauses = t;
+ }
+ }
+
if (global_bindings_p ())
{
char buf[128];
@@ -32375,6 +32415,16 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
expand_or_defer_fn (finish_function (0));
obstack_free (&declarator_obstack, p);
}
+ else
+ {
+ stmt = make_node (OACC_DECLARE);
+ TREE_TYPE (stmt) = void_type_node;
+ OACC_DECLARE_CLAUSES (stmt) = clauses;
+ OACC_DECLARE_RETURN_CLAUSES (stmt) = ret_clauses;
+ SET_EXPR_LOCATION (stmt, pragma_tok->location);
+
+ add_stmt (stmt);
+ }
return NULL_TREE;
}
diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c
index 056b2c1..8ace93c 100644
--- a/gcc/cp/pt.c
+++ b/gcc/cp/pt.c
@@ -13907,6 +13907,7 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
integral_constant_expression_p)
tree stmt, tmp;
+tree s;
tree r;
location_t loc;
@@ -14396,8 +14397,18 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
add_stmt (t);
break;
- case OMP_TARGET_UPDATE:
case OACC_DECLARE:
+ t = copy_node (t);
+ tmp = tsubst_omp_clauses (OACC_DECLARE_CLAUSES (t), false,
+ args, complain, in_decl);
+ OACC_DECLARE_CLAUSES (t) = tmp;
+ tmp = tsubst_omp_clauses (OACC_DECLARE_RETURN_CLAUSES (t), false,
+ args, complain, in_decl);
+ OACC_DECLARE_RETURN_CLAUSES (t) = tmp;
+ add_stmt (t);
+ break;
+
+ case OMP_TARGET_UPDATE:
case OACC_ENTER_DATA:
case OACC_EXIT_DATA:
case OACC_UPDATE:
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index 1e87d6c..9d76db7 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -1287,6 +1287,7 @@ typedef struct gfc_oacc_declare
locus where;
bool module_var;
gfc_omp_clauses *clauses;
+ gfc_omp_clauses *return_clauses;
}
gfc_oacc_declare;
#define gfc_get_oacc_declare() XCNEW (gfc_oacc_declare)
diff --git a/gcc/fortran/trans-decl.c b/gcc/fortran/trans-decl.c
index bc54067..eee5340 100644
--- a/gcc/fortran/trans-decl.c
+++ b/gcc/fortran/trans-decl.c
@@ -5864,8 +5864,7 @@ void
finish_oacc_declare (gfc_namespace *ns, enum sym_flavor flavor)
{
gfc_code *code, *end_c, *code2;
- gfc_oacc_declare *oc;
- gfc_omp_clauses *omp_clauses = NULL, *ret_clauses = NULL;
+ gfc_oacc_declare *oc, *new_oc;
gfc_omp_namelist *n;
locus where = gfc_current_locus;
@@ -5888,204 +5887,109 @@ finish_oacc_declare (gfc_namespace *ns, enum sym_flavor flavor)
for (oc = ns->oacc_declare; oc; oc = oc->next)
{
+ gfc_omp_clauses *omp_clauses, *ret_clauses;
+
if (oc->module_var)
continue;
if (oc->clauses)
{
- if (omp_clauses)
- {
- gfc_omp_namelist *p;
-
- for (n = omp_clauses->lists[OMP_LIST_MAP]; n; n = n->next)
- p = n;
-
- p->next = oc->clauses->lists[OMP_LIST_MAP];
- }
- else
- {
- omp_clauses = oc->clauses;
- }
- }
- }
-
- while (ns->oacc_declare)
- {
- oc = ns->oacc_declare;
- ns->oacc_declare = oc->next;
- free (oc);
- }
-
- code = XCNEW (gfc_code);
- code->op = EXEC_OACC_DECLARE;
- code->loc = where;
- code->ext.omp_clauses = omp_clauses;
-
- for (n = omp_clauses->lists[OMP_LIST_MAP]; n; n = n->next)
- {
- bool ret = false;
- gfc_omp_map_op new_op;
-
- switch (n->u.map_op)
- {
- case OMP_MAP_ALLOC:
- case OMP_MAP_FORCE_ALLOC:
- new_op = OMP_MAP_FORCE_DEALLOC;
- ret = true;
- break;
-
- case OMP_MAP_DEVICE_RESIDENT:
- n->u.map_op = OMP_MAP_FORCE_ALLOC;
- new_op = OMP_MAP_FORCE_DEALLOC;
- ret = true;
- break;
-
- case OMP_MAP_FORCE_FROM:
- n->u.map_op = OMP_MAP_FORCE_ALLOC;
- new_op = OMP_MAP_FORCE_FROM;
- ret = true;
- break;
-
- case OMP_MAP_FORCE_TO:
- new_op = OMP_MAP_FORCE_DEALLOC;
- ret = true;
- break;
-
- case OMP_MAP_FORCE_TOFROM:
- n->u.map_op = OMP_MAP_FORCE_TO;
- new_op = OMP_MAP_FORCE_FROM;
- ret = true;
- break;
-
- case OMP_MAP_FROM:
- n->u.map_op = OMP_MAP_FORCE_ALLOC;
- new_op = OMP_MAP_FROM;
- ret = true;
- break;
-
- case OMP_MAP_FORCE_DEVICEPTR:
- case OMP_MAP_FORCE_PRESENT:
- case OMP_MAP_LINK:
- case OMP_MAP_TO:
- break;
-
- case OMP_MAP_TOFROM:
- n->u.map_op = OMP_MAP_TO;
- new_op = OMP_MAP_FROM;
- ret = true;
- break;
-
- default:
- gcc_unreachable ();
- break;
- }
-
- if (ret)
- {
- gfc_omp_namelist *new_n;
+ code = XCNEW (gfc_code);
+ code->op = EXEC_OACC_DECLARE;
+ code->loc = where;
- new_n = gfc_get_omp_namelist ();
- new_n->sym = n->sym;
- new_n->u.map_op = new_op;
+ ret_clauses = NULL;
+ omp_clauses = oc->clauses;
- if (!ret_clauses)
- ret_clauses = gfc_get_omp_clauses ();
+ for (n = omp_clauses->lists[OMP_LIST_MAP]; n; n = n->next)
+ {
+ bool ret = false;
+ gfc_omp_map_op new_op;
- if (ret_clauses->lists[OMP_LIST_MAP])
- new_n->next = ret_clauses->lists[OMP_LIST_MAP];
+ switch (n->u.map_op)
+ {
+ case OMP_MAP_ALLOC:
+ case OMP_MAP_FORCE_ALLOC:
+ new_op = OMP_MAP_FORCE_DEALLOC;
+ ret = true;
+ break;
- ret_clauses->lists[OMP_LIST_MAP] = new_n;
- ret = false;
- }
- }
+ case OMP_MAP_DEVICE_RESIDENT:
+ n->u.map_op = OMP_MAP_FORCE_ALLOC;
+ new_op = OMP_MAP_FORCE_DEALLOC;
+ ret = true;
+ break;
- if (!ret_clauses)
- {
- code->next = ns->code;
- ns->code = code;
- return;
- }
+ case OMP_MAP_FORCE_FROM:
+ n->u.map_op = OMP_MAP_FORCE_ALLOC;
+ new_op = OMP_MAP_FORCE_FROM;
+ ret = true;
+ break;
- code2 = XCNEW (gfc_code);
- code2->op = EXEC_OACC_DECLARE;
- code2->loc = where;
- code2->ext.omp_clauses = ret_clauses;
+ case OMP_MAP_FORCE_TO:
+ new_op = OMP_MAP_FORCE_DEALLOC;
+ ret = true;
+ break;
- if (ns->code)
- {
- find_oacc_return (ns->code);
+ case OMP_MAP_FORCE_TOFROM:
+ n->u.map_op = OMP_MAP_FORCE_TO;
+ new_op = OMP_MAP_FORCE_FROM;
+ ret = true;
+ break;
- if (ns->code->op == EXEC_END_PROCEDURE)
- {
- code2->next = ns->code;
- code->next = code2;
- }
- else
- {
- end_c = find_end (ns->code);
- if (end_c)
- {
- code2->next = end_c->next;
- end_c->next = code2;
- code->next = ns->code;
- }
- else
- {
- gfc_code *last;
+ case OMP_MAP_FROM:
+ n->u.map_op = OMP_MAP_FORCE_ALLOC;
+ new_op = OMP_MAP_FROM;
+ ret = true;
+ break;
- last = ns->code;
+ case OMP_MAP_FORCE_DEVICEPTR:
+ case OMP_MAP_FORCE_PRESENT:
+ case OMP_MAP_LINK:
+ case OMP_MAP_TO:
+ break;
- while (last->next)
- last = last->next;
+ case OMP_MAP_TOFROM:
+ n->u.map_op = OMP_MAP_TO;
+ new_op = OMP_MAP_FROM;
+ ret = true;
+ break;
- last->next = code2;
- code->next = ns->code;
- }
- }
- }
- else
- {
- code->next = code2;
- }
+ default:
+ gcc_unreachable ();
+ break;
+ }
- while (oacc_returns)
- {
- struct oacc_return *r;
+ if (ret)
+ {
+ gfc_omp_namelist *new_n;
- r = oacc_returns;
+ new_n = gfc_get_omp_namelist ();
+ new_n->sym = n->sym;
+ new_n->u.map_op = new_op;
- ret_clauses = gfc_get_omp_clauses ();
+ if (!ret_clauses)
+ ret_clauses = gfc_get_omp_clauses ();
- for (n = omp_clauses->lists[OMP_LIST_MAP]; n; n = n->next)
- {
- if (n->u.map_op == OMP_MAP_FORCE_ALLOC
- || n->u.map_op == OMP_MAP_FORCE_TO)
- {
- gfc_omp_namelist *new_n;
+ if (ret_clauses->lists[OMP_LIST_MAP])
+ new_n->next = ret_clauses->lists[OMP_LIST_MAP];
- new_n = gfc_get_omp_namelist ();
- new_n->sym = n->sym;
- new_n->u.map_op = OMP_MAP_FORCE_DEALLOC;
+ ret_clauses->lists[OMP_LIST_MAP] = new_n;
+ ret = false;
+ }
+ }
- if (ret_clauses->lists[OMP_LIST_MAP])
- new_n->next = ret_clauses->lists[OMP_LIST_MAP];
+ code->ext.oacc_declare = gfc_get_oacc_declare ();
+ code->ext.oacc_declare->clauses = omp_clauses;
+ code->ext.oacc_declare->return_clauses = ret_clauses;
- ret_clauses->lists[OMP_LIST_MAP] = new_n;
- }
+ if (ns->code)
+ code->next = ns->code;
+ ns->code = code;
}
-
- code2 = XCNEW (gfc_code);
- code2->op = EXEC_OACC_DECLARE;
- code2->loc = where;
- code2->ext.omp_clauses = ret_clauses;
- code2->next = r->code->next;
- r->code->next = code2;
-
- oacc_returns = r->next;
- free (r);
}
- ns->code = code;
+ return;
}
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index a6b8928..cd76f2a 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -4448,7 +4448,7 @@ tree
gfc_trans_oacc_declare (gfc_code *code)
{
stmtblock_t block;
- tree stmt, oacc_clauses;
+ tree stmt, c1, c2;
enum tree_code construct_code;
gfc_start_block (&block);
@@ -4456,11 +4456,15 @@ gfc_trans_oacc_declare (gfc_code *code)
construct_code = OACC_DECLARE;
gfc_start_block (&block);
- oacc_clauses = gfc_trans_omp_clauses (&block, code->ext.omp_clauses,
- code->loc);
- stmt = build1_loc (input_location, construct_code, void_type_node,
- oacc_clauses);
+ c1 = gfc_trans_omp_clauses (&block, code->ext.oacc_declare->clauses,
+ code->loc);
+
+ c2 = gfc_trans_omp_clauses (&block, code->ext.oacc_declare->return_clauses,
+ code->loc);
+
+ stmt = build2_loc (input_location, construct_code, void_type_node, c1, c2);
gfc_add_expr_to_block (&block, stmt);
+
return gfc_finish_block (&block);
}
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index e130af9..b6d4a42 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -132,7 +132,8 @@ enum acc_region_kind
{
ARK_GENERAL, /* Default used for data, etc. regions. */
ARK_PARALLEL, /* Parallel construct. */
- ARK_KERNELS /* Kernels construct. */
+ ARK_KERNELS, /* Kernels construct. */
+ ARK_DECLARE /* Declare directive. */
};
/* Gimplify hashtable helper. */
@@ -176,6 +177,7 @@ struct gimplify_omp_ctx
enum acc_region_kind acc_region_kind;
bool combined_loop;
bool distribute;
+ gomp_target *stmt;
};
static struct gimplify_ctx *gimplify_ctxp;
@@ -7105,6 +7107,61 @@ gimplify_oacc_cache (tree *expr_p, gimple_seq *pre_p)
*expr_p = NULL_TREE;
}
+/* Gimplify OACC_DECLARE. */
+
+static void
+gimplify_oacc_declare (tree *expr_p, gimple_seq *pre_p)
+{
+ tree expr = *expr_p;
+ gomp_target *stmt;
+ tree clauses, t;
+
+ clauses = OACC_DECLARE_CLAUSES (expr);
+
+ gimplify_scan_omp_clauses (&clauses, pre_p, ORT_TARGET_DATA, ORK_OACC);
+
+ gimplify_omp_ctxp->acc_region_kind = ARK_DECLARE;
+ gimplify_omp_ctxp->stmt = NULL;
+
+ for (t = clauses; t; t = OMP_CLAUSE_CHAIN (t))
+ {
+ tree attrs, decl = OMP_CLAUSE_DECL (t);
+
+ if (TREE_CODE (decl) == MEM_REF)
+ continue;
+
+ omp_add_variable (gimplify_omp_ctxp, decl, GOVD_SEEN);
+
+ attrs = lookup_attribute ("oacc declare", DECL_ATTRIBUTES (decl));
+ if (attrs)
+ DECL_ATTRIBUTES (decl) = remove_attribute ("oacc declare", attrs);
+ }
+
+ stmt = gimple_build_omp_target (NULL, GF_OMP_TARGET_KIND_OACC_DECLARE,
+ clauses);
+
+ gimplify_seq_add_stmt (pre_p, stmt);
+
+ clauses = OACC_DECLARE_RETURN_CLAUSES (expr);
+
+ if (clauses)
+ {
+ struct gimplify_omp_ctx *c;
+
+ gimplify_scan_omp_clauses (&clauses, pre_p, ORT_TARGET_DATA, ORK_OACC);
+
+ c = gimplify_omp_ctxp;
+ gimplify_omp_ctxp = c->outer_context;
+ delete_omp_context (c);
+
+ stmt = gimple_build_omp_target (NULL, GF_OMP_TARGET_KIND_OACC_DECLARE,
+ clauses);
+ gimplify_omp_ctxp->stmt = stmt;
+ }
+
+ *expr_p = NULL_TREE;
+}
+
static tree
gimplify_oacc_host_data_1 (tree *tp, int *walk_subtrees, void *data ATTRIBUTE_UNUSED)
{
@@ -7933,10 +7990,6 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
switch (TREE_CODE (expr))
{
- case OACC_DECLARE:
- kind = GF_OMP_TARGET_KIND_OACC_DECLARE;
- ork = ORK_OACC;
- break;
case OACC_ENTER_DATA:
kind = GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA;
ork = ORK_OACC;
@@ -8914,6 +8967,11 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
case OACC_HOST_DATA:
ret = gimplify_oacc_host_data (expr_p, pre_p);
break;
+
+ case OACC_DECLARE:
+ gimplify_oacc_declare (expr_p, pre_p);
+ ret = GS_ALL_DONE;
+ break;
case OACC_KERNELS:
case OACC_PARALLEL:
@@ -8927,7 +8985,6 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
ret = GS_ALL_DONE;
break;
- case OACC_DECLARE:
case OACC_ENTER_DATA:
case OACC_EXIT_DATA:
case OACC_UPDATE:
@@ -9568,6 +9625,25 @@ gimplify_body (tree fndecl, bool do_parms)
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;
+ }
+
+ 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
not the case, wrap everything in a GIMPLE_BIND to make it so. */
if (gimple_code (outer_stmt) == GIMPLE_BIND
diff --git a/gcc/testsuite/gfortran.dg/goacc/declare-1.f95 b/gcc/testsuite/gfortran.dg/goacc/declare-1.f95
index 7ea58ef..3129f04 100644
--- a/gcc/testsuite/gfortran.dg/goacc/declare-1.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/declare-1.f95
@@ -16,4 +16,3 @@ contains
end function foo
end program test
! { dg-final { scan-tree-dump-times "pragma acc declare map\\(force_to:i\\)" 2 "original" } }
-! { dg-final { scan-tree-dump-times "pragma acc declare map\\(force_from:i\\)" 2 "original" } }
diff --git a/gcc/tree.def b/gcc/tree.def
index 56580af..9ea537e 100644
--- a/gcc/tree.def
+++ b/gcc/tree.def
@@ -1056,6 +1056,11 @@ DEFTREECODE (OACC_KERNELS, "oacc_kernels", tcc_statement, 2)
DEFTREECODE (OACC_DATA, "oacc_data", tcc_statement, 2)
+/* OpenACC - #pragma acc declare [clause1 ... clauseN]
+ Operand 0: OACC_DECLARE_CLAUSES: List of clauses.
+ Operand 1: OACC_DECLARE_RETURN_CLAUSES: List of clauses for returns. */
+DEFTREECODE (OACC_DECLARE, "oacc_declare", tcc_statement, 2)
+
/* OpenACC - #pragma acc host_data [clause1 ... clauseN]
Operand 0: OACC_HOST_DATA_BODY: Host_data construct body.
Operand 1: OACC_HOST_DATA_CLAUSES: List of clauses. */
@@ -1166,10 +1171,6 @@ DEFTREECODE (OMP_CRITICAL, "omp_critical", tcc_statement, 2)
OMP_CLAUSE__CACHE_ clauses). */
DEFTREECODE (OACC_CACHE, "oacc_cache", tcc_statement, 1)
-/* OpenACC - #pragma acc declare [clause1 ... clauseN]
- Operand 0: OACC_DECLARE_CLAUSES: List of clauses. */
-DEFTREECODE (OACC_DECLARE, "oacc_declare", tcc_statement, 1)
-
/* OpenACC - #pragma acc enter data [clause1 ... clauseN]
Operand 0: OACC_ENTER_DATA_CLAUSES: List of clauses. */
DEFTREECODE (OACC_ENTER_DATA, "oacc_enter_data", tcc_statement, 1)
diff --git a/gcc/tree.h b/gcc/tree.h
index 29bce01..5b2e267 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1233,6 +1233,8 @@ extern void protected_set_expr_location (tree, location_t);
#define OACC_DECLARE_CLAUSES(NODE) \
TREE_OPERAND (OACC_DECLARE_CHECK (NODE), 0)
+#define OACC_DECLARE_RETURN_CLAUSES(NODE) \
+ TREE_OPERAND (OACC_DECLARE_CHECK (NODE), 1)
#define OACC_ENTER_DATA_CLAUSES(NODE) \
TREE_OPERAND (OACC_ENTER_DATA_CHECK (NODE), 0)
diff --git a/libgomp/testsuite/libgomp.oacc-c++/declare-1.C b/libgomp/testsuite/libgomp.oacc-c++/declare-1.C
index 6618b10..e82a8e5 100644
--- a/libgomp/testsuite/libgomp.oacc-c++/declare-1.C
+++ b/libgomp/testsuite/libgomp.oacc-c++/declare-1.C
@@ -1,3 +1,4 @@
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
template<class T>
T foo()
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/default-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/default-1.c
index 0d75f49..3dfde71 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/default-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/default-1.c
@@ -6,6 +6,8 @@ int
main (int argc, char **argv)
{
float a, b;
+ float c;
+#pragma acc declare create (c)
a = 2.0;
b = 0.0;
@@ -60,5 +62,28 @@ main (int argc, char **argv)
if (a != 5.0)
abort ();
+#pragma acc parallel default (none) copy (a)
+ {
+ c = a;
+ a = 1.0;
+ a = a + c;
+ }
+
+ if (a != 6.0)
+ abort ();
+
+#pragma acc data copy (a)
+ {
+#pragma acc parallel default (none)
+ {
+ c = a;
+ a = 1.0;
+ a = a + c;
+ }
+ }
+
+ if (a != 7.0)
+ abort ();
+
return 0;
}
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-1.f90
index 4d58e70..18dd1bb 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/declare-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-1.f90
@@ -104,13 +104,11 @@ subroutine subr2 (a, b, c)
end subroutine
-subroutine subr1 (a, b, c)
+subroutine subr1 (a)
integer, parameter :: N = 8
integer :: i
integer :: a(N)
!$acc declare present (a)
- integer :: b(N)
- integer :: c(N)
i = 0
@@ -144,7 +142,7 @@ subroutine subr0 (a, b, c, d)
call test (b, .false.)
call test (c, .false.)
- call subr1 (a, b, c)
+ call subr1 (a)
call test (a, .true.)
call test (b, .false.)
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/default-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/default-1.f90
index f82316e..1059089 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/default-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/default-1.f90
@@ -3,6 +3,8 @@
program main
implicit none
real a, b
+ real c
+ !$acc declare create (c)
a = 2.0
b = 0.0
@@ -31,4 +33,22 @@ program main
if (a .ne. 5.0) call abort
+ !$acc parallel default (none) copy (a)
+ c = a
+ a = 1.0
+ a = a + c
+ !$acc end parallel
+
+ if (a .ne. 6.0) call abort
+
+ !$acc data copy (a)
+ !$acc parallel default (none)
+ c = a
+ a = 1.0
+ a = a + c
+ !$acc end parallel
+ !$acc end data
+
+ if (a .ne. 7.0) call abort
+
end program main
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [gomp4] OpenACC first private
2015-08-10 19:50 ` Thomas Schwinge
@ 2015-08-18 23:30 ` Thomas Schwinge
2015-08-19 12:41 ` Nathan Sidwell
0 siblings, 1 reply; 8+ messages in thread
From: Thomas Schwinge @ 2015-08-18 23:30 UTC (permalink / raw)
To: Nathan Sidwell; +Cc: james norris, GCC Patches
[-- Attachment #1: Type: text/plain, Size: 6760 bytes --]
Hi!
On Mon, 10 Aug 2015 21:50:21 +0200, I wrote:
> On Mon, 3 Aug 2015 10:30:49 -0400, Nathan Sidwell <nathan@acm.org> wrote:
> > 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.
>
> I'm seeing the following regressions after this r226508 commit -- are
> those the ones that Jim is working on resolving?
With Jim's recent commit to gomp-4_0-branch, r226970,
<http://news.gmane.org/find-root.php?message_id=%3C55CB9109.6030202%40codesourcery.com%3E>,
the following regressions are resolved (thanks!):
> PASS: libgomp.oacc-fortran/declare-1.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O0 (test for excess errors)
> [-PASS:-]{+FAIL:+} libgomp.oacc-fortran/declare-1.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O0 execution test
> PASS: libgomp.oacc-fortran/declare-1.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O1 (test for excess errors)
> [-PASS:-]{+FAIL:+} libgomp.oacc-fortran/declare-1.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O1 execution test
> PASS: libgomp.oacc-fortran/declare-1.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O2 (test for excess errors)
> [-PASS:-]{+FAIL:+} libgomp.oacc-fortran/declare-1.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O2 execution test
> PASS: libgomp.oacc-fortran/declare-1.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -fomit-frame-pointer (test for excess errors)
> [-PASS:-]{+FAIL:+} libgomp.oacc-fortran/declare-1.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -fomit-frame-pointer execution test
> PASS: libgomp.oacc-fortran/declare-1.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -fomit-frame-pointer -funroll-loops (test for excess errors)
> [-PASS:-]{+FAIL:+} libgomp.oacc-fortran/declare-1.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -fomit-frame-pointer -funroll-loops execution test
> PASS: libgomp.oacc-fortran/declare-1.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -fomit-frame-pointer -funroll-all-loops -finline-functions (test for excess errors)
> [-PASS:-]{+FAIL:+} libgomp.oacc-fortran/declare-1.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -fomit-frame-pointer -funroll-all-loops -finline-functions execution test
> PASS: libgomp.oacc-fortran/declare-1.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -g (test for excess errors)
> [-PASS:-]{+FAIL:+} libgomp.oacc-fortran/declare-1.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -g execution test
> PASS: libgomp.oacc-fortran/declare-1.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -Os (test for excess errors)
> [-PASS:-]{+FAIL:+} libgomp.oacc-fortran/declare-1.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -Os execution test
..., but the following ones remain to be addressed -- could somebody look
into this, please? Especially the timeouts are very annoying. Tests
that now reproducibly XPASS instead of XFAIL should be verified, and the
XFAIL marker removed.
> [-PASS:-]{+FAIL: gfortran.dg/goacc/modules.f95 -O (internal compiler error)+}
> {+FAIL:+} gfortran.dg/goacc/modules.f95 -O (test for excess errors)
>
> PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-loop-1.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 (test for excess errors)
> [-XFAIL:-]{+XPASS:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-loop-1.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 execution test
>
> PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-4.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 (test for excess errors)
> {+WARNING: program timed out.+}
> XFAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-4.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 execution test
>
> PASS: libgomp.oacc-c++/../libgomp.oacc-c-c++-common/parallel-loop-1.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 (test for excess errors)
> [-XFAIL:-]{+XPASS:+} libgomp.oacc-c++/../libgomp.oacc-c-c++-common/parallel-loop-1.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 execution test
>
> PASS: libgomp.oacc-c++/../libgomp.oacc-c-c++-common/reduction-4.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 (test for excess errors)
> {+WARNING: program timed out.+}
> XFAIL: libgomp.oacc-c++/../libgomp.oacc-c-c++-common/reduction-4.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 execution test
>
> PASS: libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O0 (test for excess errors)
> [-PASS:-]{+FAIL:+} libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O0 execution test
> PASS: libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O1 (test for excess errors)
> [-PASS:-]{+FAIL:+} libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O1 execution test
> PASS: libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O2 (test for excess errors)
> PASS: libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O2 execution test
> PASS: libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -fomit-frame-pointer (test for excess errors)
> PASS: libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -fomit-frame-pointer execution test
> PASS: libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -fomit-frame-pointer -funroll-loops (test for excess errors)
> [-PASS:-]{+FAIL:+} libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -fomit-frame-pointer -funroll-loops execution test
> PASS: libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -fomit-frame-pointer -funroll-all-loops -finline-functions (test for excess errors)
> PASS: libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -fomit-frame-pointer -funroll-all-loops -finline-functions execution test
> PASS: libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -g (test for excess errors)
Grüße,
Thomas
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [gomp4] OpenACC first private
2015-08-18 23:30 ` Thomas Schwinge
@ 2015-08-19 12:41 ` Nathan Sidwell
0 siblings, 0 replies; 8+ messages in thread
From: Nathan Sidwell @ 2015-08-19 12:41 UTC (permalink / raw)
To: Thomas Schwinge; +Cc: james norris, GCC Patches
On 08/18/15 17:43, Thomas Schwinge wrote:
> ..., but the following ones remain to be addressed -- could somebody look
> into this, please? Especially the timeouts are very annoying. Tests
> that now reproducibly XPASS instead of XFAIL should be verified, and the
> XFAIL marker removed.
>
>> [-PASS:-]{+FAIL: gfortran.dg/goacc/modules.f95 -O (internal compiler error)+}
>> {+FAIL:+} gfortran.dg/goacc/modules.f95 -O (test for excess errors)
>>
>> PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-loop-1.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 (test for excess errors)
>> [-XFAIL:-]{+XPASS:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-loop-1.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 execution test
>>
>> PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-4.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 (test for excess errors)
>> {+WARNING: program timed out.+}
>> XFAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/reduction-4.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 execution test
>>
>> PASS: libgomp.oacc-c++/../libgomp.oacc-c-c++-common/parallel-loop-1.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 (test for excess errors)
>> [-XFAIL:-]{+XPASS:+} libgomp.oacc-c++/../libgomp.oacc-c-c++-common/parallel-loop-1.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 execution test
>>
>> PASS: libgomp.oacc-c++/../libgomp.oacc-c-c++-common/reduction-4.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 (test for excess errors)
>> {+WARNING: program timed out.+}
>> XFAIL: libgomp.oacc-c++/../libgomp.oacc-c-c++-common/reduction-4.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 execution test
>>
>> PASS: libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O0 (test for excess errors)
>> [-PASS:-]{+FAIL:+} libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O0 execution test
>> PASS: libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O1 (test for excess errors)
>> [-PASS:-]{+FAIL:+} libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O1 execution test
>> PASS: libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O2 (test for excess errors)
>> PASS: libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O2 execution test
>> PASS: libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -fomit-frame-pointer (test for excess errors)
>> PASS: libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -fomit-frame-pointer execution test
>> PASS: libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -fomit-frame-pointer -funroll-loops (test for excess errors)
>> [-PASS:-]{+FAIL:+} libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -fomit-frame-pointer -funroll-loops execution test
>> PASS: libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -fomit-frame-pointer -funroll-all-loops -finline-functions (test for excess errors)
>> PASS: libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -fomit-frame-pointer -funroll-all-loops -finline-functions execution test
>> PASS: libgomp.oacc-fortran/lib-13.f90 -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -O3 -g (test for excess errors)
>
If the reduction ones are timeing out, they should simply be skipped until the
reduction reworking is complete. I do not know what the lib-13 ones are.
--
Nathan Sidwell
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [gomp4] declare directive
2015-08-12 18:31 ` [gomp4] declare directive James Norris
@ 2015-09-01 16:34 ` Tom de Vries
2015-09-02 8:25 ` Tom de Vries
1 sibling, 0 replies; 8+ messages in thread
From: Tom de Vries @ 2015-09-01 16:34 UTC (permalink / raw)
To: James Norris, gcc-patches; +Cc: Jakub Jelinek
[-- Attachment #1: Type: text/plain, Size: 427 bytes --]
On 12/08/15 20:31, James Norris wrote:
> diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c
> index 056b2c1..8ace93c 100644
> --- a/gcc/cp/pt.c
> +++ b/gcc/cp/pt.c
> @@ -13907,6 +13907,7 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
> integral_constant_expression_p)
>
> tree stmt, tmp;
> +tree s;
This caused a bootstrap failure (unused variable 's'). Committed fix as
attached.
Thanks,
- Tom
[-- Attachment #2: 0004-Remove-unused-variable-in-tsubst_expr.patch --]
[-- Type: text/x-patch, Size: 507 bytes --]
Remove unused variable in tsubst_expr
2015-09-01 Tom de Vries <tom@codesourcery.com>
* pt.c (tsubst_expr): Remove unused variable s.
---
gcc/cp/pt.c | 1 -
1 file changed, 1 deletion(-)
diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c
index bcea026..c94c463 100644
--- a/gcc/cp/pt.c
+++ b/gcc/cp/pt.c
@@ -14418,7 +14418,6 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
integral_constant_expression_p)
tree stmt, tmp;
-tree s;
tree r;
location_t loc;
--
1.9.1
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [gomp4] declare directive
2015-08-12 18:31 ` [gomp4] declare directive James Norris
2015-09-01 16:34 ` Tom de Vries
@ 2015-09-02 8:25 ` Tom de Vries
1 sibling, 0 replies; 8+ messages in thread
From: Tom de Vries @ 2015-09-02 8:25 UTC (permalink / raw)
To: James Norris, gcc-patches; +Cc: Jakub Jelinek
[-- Attachment #1: Type: text/plain, Size: 695 bytes --]
On 12-08-15 20:31, James Norris wrote:
> diff --git a/gcc/fortran/trans-decl.c b/gcc/fortran/trans-decl.c
> index bc54067..eee5340 100644
> --- a/gcc/fortran/trans-decl.c
> +++ b/gcc/fortran/trans-decl.c
> @@ -5864,8 +5864,7 @@ void
> finish_oacc_declare (gfc_namespace *ns, enum sym_flavor flavor)
> {
> gfc_code *code, *end_c, *code2;
> - gfc_oacc_declare *oc;
> - gfc_omp_clauses *omp_clauses = NULL, *ret_clauses = NULL;
> + gfc_oacc_declare *oc, *new_oc;
> gfc_omp_namelist *n;
> locus where = gfc_current_locus;
>
This introduces an unused variable new_oc.
Attached patch removes that and some other unused variables in finish_oacc_declare.
Committed.
Thanks,
- Tom
[-- Attachment #2: 0001-Remove-unused-vars-in-finish_oacc_declare.patch --]
[-- Type: text/x-patch, Size: 736 bytes --]
Remove unused vars in finish_oacc_declare
2015-09-02 Tom de Vries <tom@codesourcery.com>
* trans-decl.c (finish_oacc_declare): Remove unused variables.
---
gcc/fortran/trans-decl.c | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/gcc/fortran/trans-decl.c b/gcc/fortran/trans-decl.c
index c84b098..39acabd 100644
--- a/gcc/fortran/trans-decl.c
+++ b/gcc/fortran/trans-decl.c
@@ -5890,8 +5890,8 @@ find_module_oacc_declare_clauses (gfc_symbol *sym)
void
finish_oacc_declare (gfc_namespace *ns, enum sym_flavor flavor)
{
- gfc_code *code, *end_c, *code2;
- gfc_oacc_declare *oc, *new_oc;
+ gfc_code *code;
+ gfc_oacc_declare *oc;
gfc_omp_namelist *n;
locus where = gfc_current_locus;
--
1.9.1
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [gomp4] OpenACC first private
2015-08-03 14:30 [gomp4] OpenACC first private Nathan Sidwell
2015-08-10 19:50 ` Thomas Schwinge
2015-08-12 18:31 ` [gomp4] declare directive James Norris
@ 2015-10-29 8:31 ` Thomas Schwinge
2 siblings, 0 replies; 8+ messages in thread
From: Thomas Schwinge @ 2015-10-29 8:31 UTC (permalink / raw)
To: GCC Patches; +Cc: Nathan Sidwell
[-- Attachment #1: Type: text/plain, Size: 3112 bytes --]
Hi!
On Mon, 3 Aug 2015 10:30:49 -0400, Nathan Sidwell <nathan@acm.org> wrote:
> 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.
> --- 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))
Turns out, this has been the last (only?) user of
gfc_match_omp_map_clause to specify false for "allow_sections". We once
had added the latter; removed on gomp-4_0-branch in r229516:
commit 64fec7e145a784ec1e5844a8296e8a39aeea092d
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Thu Oct 29 08:21:11 2015 +0000
Cleanup: gfc_match_omp_map_clause
gcc/fortran/
* openmp.c (gfc_match_omp_map_clause): Remove allow_sections
formal parameter.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@229516 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/fortran/ChangeLog.gomp | 5 +++++
gcc/fortran/openmp.c | 8 +++-----
2 files changed, 8 insertions(+), 5 deletions(-)
diff --git gcc/fortran/ChangeLog.gomp gcc/fortran/ChangeLog.gomp
index 2cc161a..7fe3eac 100644
--- gcc/fortran/ChangeLog.gomp
+++ gcc/fortran/ChangeLog.gomp
@@ -1,3 +1,8 @@
+2015-10-29 Thomas Schwinge <thomas@codesourcery.com>
+
+ * openmp.c (gfc_match_omp_map_clause): Remove allow_sections
+ formal parameter.
+
2015-10-28 Cesar Philippidis <cesar@codesourcery.com>
* trans-openmp.c (gfc_filter_oacc_combined_clauses): Don't zero-
diff --git gcc/fortran/openmp.c gcc/fortran/openmp.c
index afcce9a..a2c5105 100644
--- gcc/fortran/openmp.c
+++ gcc/fortran/openmp.c
@@ -482,12 +482,10 @@ match_oacc_clause_gang (gfc_omp_clauses *cp)
mapping. */
static bool
-gfc_match_omp_map_clause (gfc_omp_namelist **list, gfc_omp_map_op map_op,
- bool allow_sections = true)
+gfc_match_omp_map_clause (gfc_omp_namelist **list, gfc_omp_map_op map_op)
{
gfc_omp_namelist **head = NULL;
- if (gfc_match_omp_variable_list ("", list, false, NULL, &head,
- allow_sections)
+ if (gfc_match_omp_variable_list ("", list, false, NULL, &head, true)
== MATCH_YES)
{
gfc_omp_namelist *n;
@@ -592,7 +590,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
&& gfc_match_omp_variable_list ("firstprivate (",
&c->lists[OMP_LIST_FIRSTPRIVATE],
true)
- == MATCH_YES)
+ == MATCH_YES)
continue;
if ((mask & OMP_CLAUSE_LASTPRIVATE)
&& gfc_match_omp_variable_list ("lastprivate (",
Grüße
Thomas
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]
^ permalink raw reply [flat|nested] 8+ messages in thread
end of thread, other threads:[~2015-10-29 8:23 UTC | newest]
Thread overview: 8+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2015-08-03 14:30 [gomp4] OpenACC first private Nathan Sidwell
2015-08-10 19:50 ` Thomas Schwinge
2015-08-18 23:30 ` Thomas Schwinge
2015-08-19 12:41 ` Nathan Sidwell
2015-08-12 18:31 ` [gomp4] declare directive James Norris
2015-09-01 16:34 ` Tom de Vries
2015-09-02 8:25 ` Tom de Vries
2015-10-29 8:31 ` [gomp4] OpenACC first private Thomas Schwinge
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).