From 182190f2b3a236eb30bd717981e1f2c9b51cea06 Mon Sep 17 00:00:00 2001 From: Nathan Sidwell Date: Wed, 11 Nov 2015 14:24:09 +0000 Subject: [PATCH] gimplify.c (enum omp_region_type): Add ORT_ACC, ORT_ACC_DATA, ORT_ACC_PARALLEL, ORT_ACC_KERNELS. gcc/ * gcc/gimplify.c (enum omp_region_type): Add ORT_ACC, ORT_ACC_DATA, ORT_ACC_PARALLEL, ORT_ACC_KERNELS. Adjust ORT_NONE. (gimple_add_tmp_var): Add ORT_ACC checks. (gimplify_var_or_parm_decl): Likewise. (omp_firstprivatize_variable): Likewise. Use ORT_TARGET_DATA as a mask. (omp_add_variable): Look in outer contexts for openacc and allow reductions with other sharing. Add ORT_ACC and ORT_TARGET_DATA checks. (omp_notice_variable, omp_is_private, omp_check_private): Add ORT_ACC checks. (gimplify_scan_omp_clauses: Treat ORT_ACC as ORT_WORKSHARE. Permit private openacc reductions. (gimplify_oacc_cache): Specify ORT_ACC. (gimplify_omp_workshare): Adjust OpenACC region types. (gimplify_omp_target_update): Likewise. * gcc/omp-low.c (scan_sharing_clauses): Remove Openacc firstprivate sorry. (lower-rec_input_clauses): Don't handle openacc firstprivate references here. (lower_omp_target): Emit initializers for openacc firstprivate vars. gcc/testsuite/ * gfortran.dg/goacc/private-3.f95: Remove xfail. * gfortran.dg/goacc/combined_loop.f90: Remove xfail. libgomp/ * testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c: Remove xfail. * testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c: Remove xfail. * testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c: New. * testsuite/libgomp.oacc-c-c++-common/firstprivate-2.c: New. Co-Authored-By: Cesar Philippidis From-SVN: r230169 --- gcc/ChangeLog | 26 ++++ gcc/gimplify.c | 141 ++++++++++++------ gcc/omp-low.c | 93 +++++++++--- gcc/testsuite/ChangeLog | 5 + .../gfortran.dg/goacc/combined_loop.f90 | 2 - gcc/testsuite/gfortran.dg/goacc/private-3.f95 | 2 - libgomp/ChangeLog | 5 + .../firstprivate-1.c | 41 +++++ .../firstprivate-2.c | 31 ++++ .../libgomp.oacc-c-c++-common/loop-red-v-2.c | 2 - .../libgomp.oacc-c-c++-common/loop-red-w-2.c | 2 - 11 files changed, 278 insertions(+), 72 deletions(-) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-2.c diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 57329ed9ff1..b3bbc21c87a 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,29 @@ +2015-11-11 Nathan Sidwell + Cesar Philippidis + + gcc/ + * gcc/gimplify.c (enum omp_region_type): Add ORT_ACC, + ORT_ACC_DATA, ORT_ACC_PARALLEL, ORT_ACC_KERNELS. Adjust ORT_NONE. + (gimple_add_tmp_var): Add ORT_ACC checks. + (gimplify_var_or_parm_decl): Likewise. + (omp_firstprivatize_variable): Likewise. Use ORT_TARGET_DATA as a + mask. + (omp_add_variable): Look in outer contexts for openacc and allow + reductions with other sharing. Add ORT_ACC and ORT_TARGET_DATA + checks. + (omp_notice_variable, omp_is_private, omp_check_private): Add + ORT_ACC checks. + (gimplify_scan_omp_clauses: Treat ORT_ACC as ORT_WORKSHARE. + Permit private openacc reductions. + (gimplify_oacc_cache): Specify ORT_ACC. + (gimplify_omp_workshare): Adjust OpenACC region types. + (gimplify_omp_target_update): Likewise. + * gcc/omp-low.c (scan_sharing_clauses): Remove Openacc + firstprivate sorry. + (lower-rec_input_clauses): Don't handle openacc firstprivate + references here. + (lower_omp_target): Emit initializers for openacc firstprivate vars. + 2015-11-11 Eric Botcazou PR target/67265 diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 287e51e2016..66e5168746f 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -95,22 +95,34 @@ enum gimplify_omp_var_data enum omp_region_type { - ORT_WORKSHARE = 0, - ORT_SIMD = 1, - ORT_PARALLEL = 2, - ORT_COMBINED_PARALLEL = 3, - ORT_TASK = 4, - ORT_UNTIED_TASK = 5, - ORT_TEAMS = 8, - ORT_COMBINED_TEAMS = 9, + ORT_WORKSHARE = 0x00, + ORT_SIMD = 0x01, + + ORT_PARALLEL = 0x02, + ORT_COMBINED_PARALLEL = 0x03, + + ORT_TASK = 0x04, + ORT_UNTIED_TASK = 0x05, + + ORT_TEAMS = 0x08, + ORT_COMBINED_TEAMS = 0x09, + /* Data region. */ - ORT_TARGET_DATA = 16, + ORT_TARGET_DATA = 0x10, + /* Data region with offloading. */ - ORT_TARGET = 32, - ORT_COMBINED_TARGET = 33, + ORT_TARGET = 0x20, + ORT_COMBINED_TARGET = 0x21, + + /* OpenACC variants. */ + ORT_ACC = 0x40, /* A generic OpenACC region. */ + ORT_ACC_DATA = ORT_ACC | ORT_TARGET_DATA, /* Data construct. */ + ORT_ACC_PARALLEL = ORT_ACC | ORT_TARGET, /* Parallel construct */ + ORT_ACC_KERNELS = ORT_ACC | ORT_TARGET | 0x80, /* Kernels construct. */ + /* Dummy OpenMP region, used to disable expansion of DECL_VALUE_EXPRs in taskloop pre body. */ - ORT_NONE = 64 + ORT_NONE = 0x100 }; /* Gimplify hashtable helper. */ @@ -689,7 +701,8 @@ gimple_add_tmp_var (tree tmp) struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp; while (ctx && (ctx->region_type == ORT_WORKSHARE - || ctx->region_type == ORT_SIMD)) + || ctx->region_type == ORT_SIMD + || ctx->region_type == ORT_ACC)) ctx = ctx->outer_context; if (ctx) omp_add_variable (ctx, tmp, GOVD_LOCAL | GOVD_SEEN); @@ -1804,7 +1817,8 @@ gimplify_var_or_parm_decl (tree *expr_p) struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp; while (ctx && (ctx->region_type == ORT_WORKSHARE - || ctx->region_type == ORT_SIMD)) + || ctx->region_type == ORT_SIMD + || ctx->region_type == ORT_ACC)) ctx = ctx->outer_context; if (!ctx && !nonlocal_vlas->add (decl)) { @@ -5579,7 +5593,8 @@ omp_firstprivatize_variable (struct gimplify_omp_ctx *ctx, tree decl) } else if (ctx->region_type != ORT_WORKSHARE && ctx->region_type != ORT_SIMD - && ctx->region_type != ORT_TARGET_DATA) + && ctx->region_type != ORT_ACC + && !(ctx->region_type & ORT_TARGET_DATA)) omp_add_variable (ctx, decl, GOVD_FIRSTPRIVATE); ctx = ctx->outer_context; @@ -5667,11 +5682,13 @@ omp_add_variable (struct gimplify_omp_ctx *ctx, tree decl, unsigned int flags) /* We shouldn't be re-adding the decl with the same data sharing class. */ gcc_assert ((n->value & GOVD_DATA_SHARE_CLASS & flags) == 0); - /* The only combination of data sharing classes we should see is - FIRSTPRIVATE and LASTPRIVATE. */ nflags = n->value | flags; - gcc_assert ((nflags & GOVD_DATA_SHARE_CLASS) - == (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE) + /* The only combination of data sharing classes we should see is + FIRSTPRIVATE and LASTPRIVATE. However, OpenACC permits + reduction variables to be used in data sharing clauses. */ + gcc_assert ((ctx->region_type & ORT_ACC) != 0 + || ((nflags & GOVD_DATA_SHARE_CLASS) + == (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE)) || (flags & GOVD_DATA_SHARE_CLASS) == 0); n->value = nflags; return; @@ -5968,20 +5985,47 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code) else if (is_scalar) nflags |= GOVD_FIRSTPRIVATE; } - tree type = TREE_TYPE (decl); - if (nflags == flags - && gimplify_omp_ctxp->target_firstprivatize_array_bases - && lang_hooks.decls.omp_privatize_by_reference (decl)) - type = TREE_TYPE (type); - if (nflags == flags - && !lang_hooks.types.omp_mappable_type (type)) + + struct gimplify_omp_ctx *octx = ctx->outer_context; + if ((ctx->region_type & ORT_ACC) && octx) { - error ("%qD referenced in target region does not have " - "a mappable type", decl); - nflags |= GOVD_MAP | GOVD_EXPLICIT; + /* Look in outer OpenACC contexts, to see if there's a + data attribute for this variable. */ + omp_notice_variable (octx, decl, in_code); + + for (; octx; octx = octx->outer_context) + { + if (!(octx->region_type & (ORT_TARGET_DATA | ORT_TARGET))) + break; + splay_tree_node n2 + = splay_tree_lookup (octx->variables, + (splay_tree_key) decl); + if (n2) + { + nflags |= GOVD_MAP; + goto found_outer; + } + } } - else if (nflags == flags) - nflags |= GOVD_MAP; + + { + tree type = TREE_TYPE (decl); + + if (nflags == flags + && gimplify_omp_ctxp->target_firstprivatize_array_bases + && lang_hooks.decls.omp_privatize_by_reference (decl)) + type = TREE_TYPE (type); + if (nflags == flags + && !lang_hooks.types.omp_mappable_type (type)) + { + error ("%qD referenced in target region does not have " + "a mappable type", decl); + nflags |= GOVD_MAP | GOVD_EXPLICIT; + } + else if (nflags == flags) + nflags |= GOVD_MAP; + } + found_outer: omp_add_variable (ctx, decl, nflags); } else @@ -5998,7 +6042,8 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code) { if (ctx->region_type == ORT_WORKSHARE || ctx->region_type == ORT_SIMD - || ctx->region_type == ORT_TARGET_DATA) + || ctx->region_type == ORT_ACC + || (ctx->region_type & ORT_TARGET_DATA) != 0) goto do_outer; flags = omp_default_clause (ctx, decl, in_code, flags); @@ -6112,7 +6157,8 @@ omp_is_private (struct gimplify_omp_ctx *ctx, tree decl, int simd) } if (ctx->region_type != ORT_WORKSHARE - && ctx->region_type != ORT_SIMD) + && ctx->region_type != ORT_SIMD + && ctx->region_type != ORT_ACC) return false; else if (ctx->outer_context) return omp_is_private (ctx->outer_context, decl, simd); @@ -6168,7 +6214,8 @@ omp_check_private (struct gimplify_omp_ctx *ctx, tree decl, bool copyprivate) } } while (ctx->region_type == ORT_WORKSHARE - || ctx->region_type == ORT_SIMD); + || ctx->region_type == ORT_SIMD + || ctx->region_type == ORT_ACC); return false; } @@ -6311,7 +6358,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, omp_notice_variable (outer_ctx->outer_context, decl, true); } else if (outer_ctx - && outer_ctx->region_type == ORT_WORKSHARE + && (outer_ctx->region_type == ORT_WORKSHARE + || outer_ctx->region_type == ORT_ACC) && outer_ctx->combined_loop && splay_tree_lookup (outer_ctx->variables, (splay_tree_key) decl) == NULL @@ -6335,7 +6383,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, goto do_add; case OMP_CLAUSE_REDUCTION: flags = GOVD_REDUCTION | GOVD_SEEN | GOVD_EXPLICIT; - check_non_private = "reduction"; + /* OpenACC permits reductions on private variables. */ + if (!(region_type & ORT_ACC)) + check_non_private = "reduction"; decl = OMP_CLAUSE_DECL (c); if (TREE_CODE (decl) == MEM_REF) { @@ -7704,7 +7754,7 @@ gimplify_oacc_cache (tree *expr_p, gimple_seq *pre_p) { tree expr = *expr_p; - gimplify_scan_omp_clauses (&OACC_CACHE_CLAUSES (expr), pre_p, ORT_WORKSHARE, + gimplify_scan_omp_clauses (&OACC_CACHE_CLAUSES (expr), pre_p, ORT_ACC, OACC_CACHE); gimplify_adjust_omp_clauses (pre_p, &OACC_CACHE_CLAUSES (expr), OACC_CACHE); @@ -7833,7 +7883,9 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p) case OMP_FOR: case CILK_FOR: case OMP_DISTRIBUTE: + break; case OACC_LOOP: + ort = ORT_ACC; break; case OMP_TASKLOOP: if (find_omp_clause (OMP_FOR_CLAUSES (for_stmt), OMP_CLAUSE_UNTIED)) @@ -8895,10 +8947,14 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p) ort = OMP_TARGET_COMBINED (expr) ? ORT_COMBINED_TARGET : ORT_TARGET; break; case OACC_KERNELS: + ort = ORT_ACC_KERNELS; + break; case OACC_PARALLEL: - ort = ORT_TARGET; + ort = ORT_ACC_PARALLEL; break; case OACC_DATA: + ort = ORT_ACC_DATA; + break; case OMP_TARGET_DATA: ort = ORT_TARGET_DATA; break; @@ -8920,7 +8976,7 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p) pop_gimplify_context (g); else pop_gimplify_context (NULL); - if (ort == ORT_TARGET_DATA) + if ((ort & ORT_TARGET_DATA) != 0) { enum built_in_function end_ix; switch (TREE_CODE (expr)) @@ -8995,17 +9051,18 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p) tree expr = *expr_p; int kind; gomp_target *stmt; + enum omp_region_type ort = ORT_WORKSHARE; switch (TREE_CODE (expr)) { case OACC_ENTER_DATA: - kind = GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA; - break; case OACC_EXIT_DATA: kind = GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA; + ort = ORT_ACC; break; case OACC_UPDATE: kind = GF_OMP_TARGET_KIND_OACC_UPDATE; + ort = ORT_ACC; break; case OMP_TARGET_UPDATE: kind = GF_OMP_TARGET_KIND_UPDATE; @@ -9020,7 +9077,7 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p) gcc_unreachable (); } gimplify_scan_omp_clauses (&OMP_STANDALONE_CLAUSES (expr), pre_p, - ORT_WORKSHARE, TREE_CODE (expr)); + ort, TREE_CODE (expr)); gimplify_adjust_omp_clauses (pre_p, &OMP_STANDALONE_CLAUSES (expr), TREE_CODE (expr)); stmt = gimple_build_omp_target (NULL, kind, OMP_STANDALONE_CLAUSES (expr)); diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 2a552da1e5b..51b471cff5a 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -1896,12 +1896,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) /* FALLTHRU */ case OMP_CLAUSE_FIRSTPRIVATE: - if (is_gimple_omp_oacc (ctx->stmt)) - { - sorry ("clause not supported yet"); - break; - } - /* FALLTHRU */ case OMP_CLAUSE_LINEAR: decl = OMP_CLAUSE_DECL (c); do_private: @@ -2167,12 +2161,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) /* FALLTHRU */ case OMP_CLAUSE_FIRSTPRIVATE: - if (is_gimple_omp_oacc (ctx->stmt)) - { - sorry ("clause not supported yet"); - break; - } - /* FALLTHRU */ case OMP_CLAUSE_PRIVATE: case OMP_CLAUSE_LINEAR: case OMP_CLAUSE_IS_DEVICE_PTR: @@ -4684,7 +4672,7 @@ lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist, gimplify_assign (ptr, x, ilist); } } - else if (is_reference (var)) + else if (is_reference (var) && !is_oacc_parallel (ctx)) { /* For references that are being privatized for Fortran, allocate new backing storage for the new pointer @@ -14911,7 +14899,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) tree child_fn, t, c; gomp_target *stmt = as_a (gsi_stmt (*gsi_p)); gbind *tgt_bind, *bind, *dep_bind = NULL; - gimple_seq tgt_body, olist, ilist, new_body; + gimple_seq tgt_body, olist, ilist, fplist, new_body; location_t loc = gimple_location (stmt); bool offloaded, data_region; unsigned int map_cnt = 0; @@ -14963,6 +14951,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) child_fn = ctx->cb.dst_fn; push_gimplify_context (); + fplist = NULL; for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c)) switch (OMP_CLAUSE_CODE (c)) @@ -15007,6 +14996,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) /* FALLTHRU */ case OMP_CLAUSE_TO: case OMP_CLAUSE_FROM: + oacc_firstprivate: var = OMP_CLAUSE_DECL (c); if (!DECL_P (var)) { @@ -15029,6 +15019,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) } if (offloaded + && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_REFERENCE)) { @@ -15057,17 +15048,40 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) x = build_receiver_ref (var, true, ctx); tree new_var = lookup_decl (var, ctx); - if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c) && TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE) x = build_simple_mem_ref (x); - SET_DECL_VALUE_EXPR (new_var, x); - DECL_HAS_VALUE_EXPR_P (new_var) = 1; + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE) + { + gcc_assert (is_gimple_omp_oacc (ctx->stmt)); + if (is_reference (new_var)) + { + /* Create a local object to hold the instance + value. */ + tree type = TREE_TYPE (TREE_TYPE (new_var)); + const char *id = IDENTIFIER_POINTER (DECL_NAME (new_var)); + tree inst = create_tmp_var (type, id); + gimplify_assign (inst, fold_indirect_ref (x), &fplist); + x = build_fold_addr_expr (inst); + } + gimplify_assign (new_var, x, &fplist); + } + else if (DECL_P (new_var)) + { + SET_DECL_VALUE_EXPR (new_var, x); + DECL_HAS_VALUE_EXPR_P (new_var) = 1; + } + else + gcc_unreachable (); } map_cnt++; break; case OMP_CLAUSE_FIRSTPRIVATE: + if (is_oacc_parallel (ctx)) + goto oacc_firstprivate; map_cnt++; var = OMP_CLAUSE_DECL (c); if (!is_reference (var) @@ -15092,6 +15106,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) break; case OMP_CLAUSE_PRIVATE: + if (is_gimple_omp_oacc (ctx->stmt)) + break; var = OMP_CLAUSE_DECL (c); if (is_variable_sized (var)) { @@ -15195,9 +15211,11 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) default: break; + case OMP_CLAUSE_MAP: case OMP_CLAUSE_TO: case OMP_CLAUSE_FROM: + oacc_firstprivate_map: nc = c; ovar = OMP_CLAUSE_DECL (c); if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP @@ -15248,9 +15266,9 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) x = build_sender_ref (ovar, ctx); if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP - && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER - && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c) - && TREE_CODE (TREE_TYPE (ovar)) == ARRAY_TYPE) + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER + && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c) + && TREE_CODE (TREE_TYPE (ovar)) == ARRAY_TYPE) { gcc_assert (offloaded); tree avar @@ -15261,6 +15279,15 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) avar = build_fold_addr_expr (avar); gimplify_assign (x, avar, &ilist); } + else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE) + { + gcc_assert (is_gimple_omp_oacc (ctx->stmt)); + if (!is_reference (var)) + var = build_fold_addr_expr (var); + else + talign = TYPE_ALIGN_UNIT (TREE_TYPE (TREE_TYPE (ovar))); + gimplify_assign (x, var, &ilist); + } else if (is_gimple_reg (var)) { gcc_assert (offloaded); @@ -15289,7 +15316,17 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) gimplify_assign (x, var, &ilist); } } - s = OMP_CLAUSE_SIZE (c); + s = NULL_TREE; + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE) + { + gcc_checking_assert (is_gimple_omp_oacc (ctx->stmt)); + s = TREE_TYPE (ovar); + if (TREE_CODE (s) == REFERENCE_TYPE) + s = TREE_TYPE (s); + s = TYPE_SIZE_UNIT (s); + } + else + s = OMP_CLAUSE_SIZE (c); if (s == NULL_TREE) s = TYPE_SIZE_UNIT (TREE_TYPE (ovar)); s = fold_convert (size_type_node, s); @@ -15330,6 +15367,11 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) tkind_zero = tkind; } break; + case OMP_CLAUSE_FIRSTPRIVATE: + gcc_checking_assert (is_gimple_omp_oacc (ctx->stmt)); + tkind = GOMP_MAP_TO; + tkind_zero = tkind; + break; case OMP_CLAUSE_TO: tkind = GOMP_MAP_TO; tkind_zero = tkind; @@ -15369,6 +15411,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) break; case OMP_CLAUSE_FIRSTPRIVATE: + if (is_oacc_parallel (ctx)) + goto oacc_firstprivate_map; ovar = OMP_CLAUSE_DECL (c); if (is_reference (ovar)) talign = TYPE_ALIGN_UNIT (TREE_TYPE (TREE_TYPE (ovar))); @@ -15543,6 +15587,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) gimple_seq_add_stmt (&new_body, gimple_build_assign (ctx->receiver_decl, t)); } + gimple_seq_add_seq (&new_body, fplist); if (offloaded || data_region) { @@ -15554,6 +15599,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) default: break; case OMP_CLAUSE_FIRSTPRIVATE: + if (is_gimple_omp_oacc (ctx->stmt)) + break; var = OMP_CLAUSE_DECL (c); if (is_reference (var) || is_gimple_reg_type (TREE_TYPE (var))) @@ -15639,6 +15686,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) } break; case OMP_CLAUSE_PRIVATE: + if (is_gimple_omp_oacc (ctx->stmt)) + break; var = OMP_CLAUSE_DECL (c); if (is_reference (var)) { @@ -15727,7 +15776,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) /* Handle GOMP_MAP_FIRSTPRIVATE_{POINTER,REFERENCE} in second pass, so that firstprivate vars holding OMP_CLAUSE_SIZE if needed are already handled. */ - for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c)) + for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) switch (OMP_CLAUSE_CODE (c)) { tree var; diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog index 78332c17682..6c2c07a006f 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,3 +1,8 @@ +2015-11-11 Nathan Sidwell + + * gfortran.dg/goacc/private-3.f95: Remove xfail. + * gfortran.dg/goacc/combined_loop.f90: Remove xfail. + 2015-11-11 Eric Botcazou * gcc.target/i386/pr67265.c: New test. diff --git a/gcc/testsuite/gfortran.dg/goacc/combined_loop.f90 b/gcc/testsuite/gfortran.dg/goacc/combined_loop.f90 index e0ea87af86d..6507ddaf33e 100644 --- a/gcc/testsuite/gfortran.dg/goacc/combined_loop.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/combined_loop.f90 @@ -1,6 +1,4 @@ ! { dg-do compile } -! -! { dg-xfail-if "TODO" { *-*-* } } ! ! PR fortran/64726 diff --git a/gcc/testsuite/gfortran.dg/goacc/private-3.f95 b/gcc/testsuite/gfortran.dg/goacc/private-3.f95 index af7d683f818..349026350d4 100644 --- a/gcc/testsuite/gfortran.dg/goacc/private-3.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/private-3.f95 @@ -1,6 +1,4 @@ ! { dg-do compile } -! -! { dg-xfail-if "TODO" { *-*-* } } ! test for private variables in a reduction clause diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index ed86943bb32..406c5720325 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,8 @@ +2015-11-1 Nathan Sidwell + + * testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c: New. + * testsuite/libgomp.oacc-c-c++-common/firstprivate-2.c: New. + 2015-11-09 Nathan Sidwell * testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c: Remove diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c new file mode 100644 index 00000000000..7f5d3d37617 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c @@ -0,0 +1,41 @@ +/* { dg-do run } */ + +#include + +int main () +{ + int ok = 1; + int val = 2; + int ary[32]; + int ondev = 0; + + for (int i = 0; i < 32; i++) + ary[i] = ~0; + +#pragma acc parallel num_gangs (32) copy (ok) firstprivate (val) copy(ary, ondev) + { + ondev = acc_on_device (acc_device_not_host); +#pragma acc loop gang(static:1) + for (unsigned i = 0; i < 32; i++) + { + if (val != 2) + ok = 0; + val += i; + ary[i] = val; + } + } + + if (ondev) + { + if (!ok) + return 1; + if (val != 2) + return 1; + + for (int i = 0; i < 32; i++) + if (ary[i] != 2 + i) + return 1; + } + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-2.c new file mode 100644 index 00000000000..9666542fd82 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-2.c @@ -0,0 +1,31 @@ +/* { dg-do run } */ + +#include + +int main () +{ + int ok = 1; + int val = 2; + +#pragma acc data copy(val) + { +#pragma acc parallel present (val) + { + val = 7; + } + +#pragma acc parallel firstprivate (val) copy(ok) + { + ok = val == 7; + val = 9; + } + + } + + if (!ok) + return 1; + if(val != 7) + return 1; + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c index fbed589e146..e66732da32c 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c @@ -1,7 +1,5 @@ /* { dg-do run } */ /* { dg-additional-options "-O2" */ -/* - { dg-xfail-if "TODO" { *-*-* } } */ #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c index 47f1da02e2e..0059077b685 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c @@ -1,7 +1,5 @@ /* { dg-do run } */ /* { dg-additional-options "-O2" */ -/* - { dg-xfail-if "TODO" { *-*-* } } */ #include -- 2.30.2