From patchwork Tue Nov 10 14:12:55 2015 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Nathan Sidwell X-Patchwork-Id: 542332 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 5E70F141415 for ; Wed, 11 Nov 2015 01:13:13 +1100 (AEDT) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=QihWe/Oj; dkim-atps=neutral DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender :subject:to:references:cc:from:message-id:date:mime-version :in-reply-to:content-type; q=dns; s=default; b=LqtATa+VCLEOFJTHk hFCpbMZnt32ToggkgveR0Fw/XmBplkqFuI55pryRXjCtwcUlvfv7PmQ+2pC674hL AWBm8lkyfeVCTPjXmd5z6IzGV0RhM5qQK/H31sG5MghVpb3UhrB2bFyApTRnuvQ0 vVVjEI25w6ZtdFCdl2RY8Hq3CI= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender :subject:to:references:cc:from:message-id:date:mime-version :in-reply-to:content-type; s=default; bh=v8evWgaOAPPjOckjQXFm7Kz NYoM=; b=QihWe/OjUtzL+WtMg909GRxWg//01XxorojEg1kBB/sHgi2V8tkDYI5 NBU8LEjgJwF5EYVPgTDUZM4I6YHkkXRvHtfQ/OWoVyfe5VLTC++WAsZ+5jiVznO2 vep/28NNo9ZeP7naAVX1gTx3YL6qIJz4BWPdrUPWyAteE6F3CZsc= Received: (qmail 1832 invoked by alias); 10 Nov 2015 14:13:03 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 1796 invoked by uid 89); 10 Nov 2015 14:13:03 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.8 required=5.0 tests=BAYES_00, FREEMAIL_FROM, KAM_ASCII_DIVIDERS, RCVD_IN_DNSWL_LOW, SPF_PASS autolearn=no version=3.3.2 X-HELO: mail-qk0-f182.google.com Received: from mail-qk0-f182.google.com (HELO mail-qk0-f182.google.com) (209.85.220.182) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES128-GCM-SHA256 encrypted) ESMTPS; Tue, 10 Nov 2015 14:12:59 +0000 Received: by qkao63 with SMTP id o63so43562690qka.2 for ; Tue, 10 Nov 2015 06:12:56 -0800 (PST) X-Received: by 10.55.207.3 with SMTP id e3mr4450487qkj.32.1447164776787; Tue, 10 Nov 2015 06:12:56 -0800 (PST) Received: from ?IPv6:2601:181:c000:c497:a2a8:cdff:fe3e:b48? ([2601:181:c000:c497:a2a8:cdff:fe3e:b48]) by smtp.googlemail.com with ESMTPSA id o77sm1365520qgd.15.2015.11.10.06.12.55 (version=TLSv1.2 cipher=ECDHE-RSA-AES128-GCM-SHA256 bits=128/128); Tue, 10 Nov 2015 06:12:56 -0800 (PST) Subject: Re: OpenACC Firstprivate To: Jakub Jelinek References: <563E01A4.20607@acm.org> <20151109134619.GQ5675@tucnak.redhat.com> <5640A6B3.3030409@acm.org> <20151109141034.GS5675@tucnak.redhat.com> <5640B1B4.2070701@acm.org> Cc: GCC Patches , Cesar Philippidis From: Nathan Sidwell Message-ID: <5641FB67.9010409@acm.org> Date: Tue, 10 Nov 2015 09:12:55 -0500 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:38.0) Gecko/20100101 Thunderbird/38.3.0 MIME-Version: 1.0 In-Reply-To: <5640B1B4.2070701@acm.org> On 11/09/15 09:46, Nathan Sidwell wrote: > I'm going to try and get clarification, but I think the intent is to initialize > with the value seen on the device. Consider: > My thinking is that the intent of the firstprivate is to initialize with the > value known on the device (and behave as-if copyin, if it's not there). Not the > value most recently seen on the host -- the update clause could change that, and > may well be being used as a debugging aide, so it seems bizarre that it can > change program semantics in such a way. We believe my example is well formed. The data clauses transfer liveness of the data from host to device (and vice versa). It is ill formed to manipulate the data on the non-live system. firstprivate's intial value is taken from the (statically determined) live location. Unless I'm misunderstanding something about GOMP_MAP_FIRSTPRIVATE, using regular target mapping is the right thing. Here's an updated patch with the other two issues you noted fixed. nathan 2015-11-10 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. 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. Index: gcc/gimplify.c =================================================================== --- gcc/gimplify.c (revision 230107) +++ gcc/gimplify.c (working copy) @@ -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 gimp } 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_ct /* 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 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)) - { - error ("%qD referenced in target region does not have " - "a mappable type", decl); - nflags |= GOVD_MAP | GOVD_EXPLICIT; + + struct gimplify_omp_ctx *octx = ctx->outer_context; + if ((ctx->region_type & ORT_ACC) && octx) + { + /* 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 { 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 } 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_c } } 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, 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, 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, gimpl { 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_s 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, gi 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, gi 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 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 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)); Index: gcc/omp-low.c =================================================================== --- gcc/omp-low.c (revision 230107) +++ gcc/omp-low.c (working copy) @@ -1896,12 +1896,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_LINEAR: decl = OMP_CLAUSE_DECL (c); do_private: @@ -2167,12 +2161,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_LINEAR: case OMP_CLAUSE_IS_DEVICE_PTR: @@ -4684,7 +4672,7 @@ lower_rec_input_clauses (tree clauses, g 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 * 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 * 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 * /* 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 * } 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 * 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 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; + } + 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 * 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 * 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 * 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 * 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 * 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 * 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 * 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 * 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 * 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 * } 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 * /* 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; Index: gcc/testsuite/gfortran.dg/goacc/private-3.f95 =================================================================== --- gcc/testsuite/gfortran.dg/goacc/private-3.f95 (revision 230107) +++ gcc/testsuite/gfortran.dg/goacc/private-3.f95 (working copy) @@ -1,6 +1,4 @@ ! { dg-do compile } -! -! { dg-xfail-if "TODO" { *-*-* } } ! test for private variables in a reduction clause Index: gcc/testsuite/gfortran.dg/goacc/combined_loop.f90 =================================================================== --- gcc/testsuite/gfortran.dg/goacc/combined_loop.f90 (revision 230107) +++ gcc/testsuite/gfortran.dg/goacc/combined_loop.f90 (working copy) @@ -1,6 +1,4 @@ ! { dg-do compile } -! -! { dg-xfail-if "TODO" { *-*-* } } ! ! PR fortran/64726 Index: libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c (revision 230107) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c (working copy) @@ -1,7 +1,5 @@ /* { dg-do run } */ /* { dg-additional-options "-O2" */ -/* - { dg-xfail-if "TODO" { *-*-* } } */ #include Index: libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c (revision 230107) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c (working copy) @@ -1,7 +1,5 @@ /* { dg-do run } */ /* { dg-additional-options "-O2" */ -/* - { dg-xfail-if "TODO" { *-*-* } } */ #include Index: libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c =================================================================== --- libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c (revision 0) +++ libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c (working copy) @@ -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; +}