2015-08-03 Nathan Sidwell * 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 (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: