From patchwork Fri Oct 23 12:12:17 2015 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Jakub Jelinek X-Patchwork-Id: 534936 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 7F51D14031F for ; Fri, 23 Oct 2015 23:12:47 +1100 (AEDT) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=G1CuX81j; 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:date :from:to:subject:message-id:reply-to:mime-version:content-type; q=dns; s=default; b=TG5zBNq/0aPUdsrOVPi611tijNBGy8w3KaW+uNkkBSn 3zhFDml+Hr/hcE5axIvVT9zGJoAS2xQSzhBfuqB8L+mul/Sq4zICM2vHB6xBktGN v7cwETRzwzF0gP+VpWyOaXy+RgDrY6ESwMg0PUKyKrJZdu7O3O8HewpPrQ0HDcvM = 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:date :from:to:subject:message-id:reply-to:mime-version:content-type; s=default; bh=vXbZQKmkD/ev/hizOX00vXW+bts=; b=G1CuX81jIXPn6Vl2G W0binVQj7lNfr0X1NIHvYiqynvU6e1GyjJ1Bq3lFzDOMWjermaI5uYB8PJ9kkNU+ jQm8gAA4yMJmA3ayABy0wgPCdmkHpD0ps/AoWvs6RqImHGAsktjk8IA2PO/crobU YHNVZvEdYULJX0rlhGpJ61tWZQ= Received: (qmail 27598 invoked by alias); 23 Oct 2015 12:12:38 -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 25048 invoked by uid 89); 23 Oct 2015 12:12:37 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.8 required=5.0 tests=AWL, BAYES_00, RP_MATCHES_RCVD, SPF_HELO_PASS autolearn=ham version=3.3.2 X-HELO: mx1.redhat.com Received: from mx1.redhat.com (HELO mx1.redhat.com) (209.132.183.28) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES256-GCM-SHA384 encrypted) ESMTPS; Fri, 23 Oct 2015 12:12:24 +0000 Received: from int-mx09.intmail.prod.int.phx2.redhat.com (int-mx09.intmail.prod.int.phx2.redhat.com [10.5.11.22]) by mx1.redhat.com (Postfix) with ESMTPS id 586808F26B for ; Fri, 23 Oct 2015 12:12:22 +0000 (UTC) Received: from tucnak.zalov.cz (ovpn-116-53.ams2.redhat.com [10.36.116.53]) by int-mx09.intmail.prod.int.phx2.redhat.com (8.14.4/8.14.4) with ESMTP id t9NCCKcQ000575 (version=TLSv1/SSLv3 cipher=DHE-RSA-AES256-GCM-SHA384 bits=256 verify=NO) for ; Fri, 23 Oct 2015 08:12:21 -0400 Received: from tucnak.zalov.cz (localhost [127.0.0.1]) by tucnak.zalov.cz (8.15.2/8.15.2) with ESMTP id t9NCCIJE022094 for ; Fri, 23 Oct 2015 14:12:19 +0200 Received: (from jakub@localhost) by tucnak.zalov.cz (8.15.2/8.15.2/Submit) id t9NCCH54022093 for gcc-patches@gcc.gnu.org; Fri, 23 Oct 2015 14:12:17 +0200 Date: Fri, 23 Oct 2015 14:12:17 +0200 From: Jakub Jelinek To: gcc-patches@gcc.gnu.org Subject: [gomp4.5] Support lastprivate on distribute Message-ID: <20151023121217.GF478@tucnak.redhat.com> Reply-To: Jakub Jelinek MIME-Version: 1.0 Content-Disposition: inline User-Agent: Mutt/1.5.23 (2014-03-12) X-IsSubscribed: yes Hi! This patch adds support for lastprivate clause on distribute construct (firstprivate+lastprivate on the same var is disallowed, as one can't synchronize between contention groups). Regtested on x86_64-linux (normal and intelmicemul offloading). 2015-10-23 Jakub Jelinek gcc/ * gimplify.c (omp_default_clause): Tweak for * private/firstprivate/is_device_ptr variables on target construct and use_device_ptr on target data. (omp_check_private): Likewise. (omp_no_lastprivate): Return true only for Fortran. (gimplify_scan_omp_clauses): Fix up handling of lastprivate and linear when combined with distribute. (gimplify_adjust_omp_clauses): Diagnose the same var on both firstprivate and lastprivate on distribute construct. (gimplify_omp_for): Fix up handling of predetermined lastprivate or linear iter vars when combined with distribute. * omp-low.c (add_taskreg_looptemp_clauses): Add one extra _looptemp_ clause even for distribute parallel for, if there are lastprivate clauses on the for. (expand_omp_for_static_nochunk, expand_omp_for_static_chunk): Initialize the extra _looptemp_ clause to fd->loop.n2. (lower_omp_for_lastprivate): Determine the right count variable for distribute simd, or distribute parallel for{, simd}. gcc/c-family/ * c-omp.c (c_omp_split_clauses): Adjust for lastprivate being allowed on distribute. gcc/c/ * c-parser.c (OMP_DISTRIBUTE_CLAUSE_MASK): Add lastprivate clause. gcc/cp/ * parser.c (OMP_DISTRIBUTE_CLAUSE_MASK): Add lastprivate clause. gcc/testsuite/ * c-c++-common/gomp/distribute-1.c: New test. * c-c++-common/gomp/pr61486-2.c: Add #pragma omp declare target and #pragma omp end declare target pair around the function. Change s from a parameter to a file scope variable. libgomp/ * testsuite/libgomp.c/pr66199-5.c: New test. * testsuite/libgomp.c/pr66199-6.c: New test. * testsuite/libgomp.c/pr66199-7.c: New test. * testsuite/libgomp.c/pr66199-8.c: New test. * testsuite/libgomp.c/pr66199-9.c: New test. * testsuite/libgomp.c++/pr66199-3.C: New test. * testsuite/libgomp.c++/pr66199-4.C: New test. * testsuite/libgomp.c++/pr66199-5.C: New test. * testsuite/libgomp.c++/pr66199-6.C: New test. * testsuite/libgomp.c++/pr66199-7.C: New test. * testsuite/libgomp.c++/pr66199-8.C: New test. * testsuite/libgomp.c++/pr66199-9.C: New test. Jakub --- gcc/gimplify.c.jj 2015-10-22 13:41:23.214882261 +0200 +++ gcc/gimplify.c 2015-10-22 18:01:45.125588093 +0200 @@ -5849,9 +5849,10 @@ omp_default_clause (struct gimplify_omp_ { splay_tree_node n2; - if ((octx->region_type & (ORT_TARGET_DATA | ORT_TARGET)) != 0) - continue; n2 = splay_tree_lookup (octx->variables, (splay_tree_key) decl); + if ((octx->region_type & (ORT_TARGET_DATA | ORT_TARGET)) != 0 + && (n2 == NULL || (n2->value & GOVD_DATA_SHARE_CLASS) == 0)) + continue; if (n2 && (n2->value & GOVD_DATA_SHARE_CLASS) != GOVD_SHARED) { flags |= GOVD_FIRSTPRIVATE; @@ -6143,10 +6144,12 @@ omp_check_private (struct gimplify_omp_c return true; } - if ((ctx->region_type & (ORT_TARGET | ORT_TARGET_DATA)) != 0) + n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl); + + if ((ctx->region_type & (ORT_TARGET | ORT_TARGET_DATA)) != 0 + && (n == NULL || (n->value & GOVD_DATA_SHARE_CLASS) == 0)) continue; - n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl); if (n != NULL) { if ((n->value & GOVD_LOCAL) != 0 @@ -6177,12 +6180,12 @@ omp_no_lastprivate (struct gimplify_omp_ if (!ctx->combined_loop) return false; if (ctx->distribute) - return true; + return lang_GNU_Fortran (); break; case ORT_COMBINED_PARALLEL: break; case ORT_COMBINED_TEAMS: - return true; + return lang_GNU_Fortran (); default: return false; } @@ -6279,16 +6282,25 @@ gimplify_scan_omp_clauses (tree *list_p, else if (error_operand_p (decl)) goto do_add; else if (outer_ctx - && outer_ctx->region_type == ORT_COMBINED_PARALLEL + && (outer_ctx->region_type == ORT_COMBINED_PARALLEL + || outer_ctx->region_type == ORT_COMBINED_TEAMS) && splay_tree_lookup (outer_ctx->variables, (splay_tree_key) decl) == NULL) - omp_add_variable (outer_ctx, decl, GOVD_SHARED | GOVD_SEEN); + { + omp_add_variable (outer_ctx, decl, GOVD_SHARED | GOVD_SEEN); + if (outer_ctx->outer_context) + omp_notice_variable (outer_ctx->outer_context, decl, true); + } else if (outer_ctx && (outer_ctx->region_type & ORT_TASK) != 0 && outer_ctx->combined_loop && splay_tree_lookup (outer_ctx->variables, (splay_tree_key) decl) == NULL) - omp_add_variable (outer_ctx, decl, GOVD_LASTPRIVATE | GOVD_SEEN); + { + omp_add_variable (outer_ctx, decl, GOVD_LASTPRIVATE | GOVD_SEEN); + if (outer_ctx->outer_context) + omp_notice_variable (outer_ctx->outer_context, decl, true); + } else if (outer_ctx && outer_ctx->region_type == ORT_WORKSHARE && outer_ctx->combined_loop @@ -6302,8 +6314,14 @@ gimplify_scan_omp_clauses (tree *list_p, == ORT_COMBINED_PARALLEL) && splay_tree_lookup (outer_ctx->outer_context->variables, (splay_tree_key) decl) == NULL) - omp_add_variable (outer_ctx->outer_context, decl, - GOVD_SHARED | GOVD_SEEN); + { + struct gimplify_omp_ctx *octx = outer_ctx->outer_context; + omp_add_variable (octx, decl, GOVD_SHARED | GOVD_SEEN); + if (octx->outer_context) + omp_notice_variable (octx->outer_context, decl, true); + } + else if (outer_ctx->outer_context) + omp_notice_variable (outer_ctx->outer_context, decl, true); } goto do_add; case OMP_CLAUSE_REDUCTION: @@ -6416,9 +6434,7 @@ gimplify_scan_omp_clauses (tree *list_p, { if (octx->outer_context && (octx->outer_context->region_type - == ORT_COMBINED_PARALLEL - || (octx->outer_context->region_type - == ORT_COMBINED_TEAMS))) + == ORT_COMBINED_PARALLEL)) octx = octx->outer_context; else if (omp_check_private (octx, decl, false)) break; @@ -6433,8 +6449,15 @@ gimplify_scan_omp_clauses (tree *list_p, && octx == outer_ctx) flags = GOVD_SEEN | GOVD_SHARED; else if (octx + && octx->region_type == ORT_COMBINED_TEAMS) + flags = GOVD_SEEN | GOVD_SHARED; + else if (octx && octx->region_type == ORT_COMBINED_TARGET) - flags &= ~GOVD_LASTPRIVATE; + { + flags &= ~GOVD_LASTPRIVATE; + if (flags == GOVD_SEEN) + break; + } else break; splay_tree_node on @@ -7289,6 +7312,15 @@ gimplify_adjust_omp_clauses (gimple_seq else OMP_CLAUSE_CODE (c) = OMP_CLAUSE_PRIVATE; } + else if (code == OMP_DISTRIBUTE + && OMP_CLAUSE_LASTPRIVATE_FIRSTPRIVATE (c)) + { + remove = true; + error_at (OMP_CLAUSE_LOCATION (c), + "same variable used in % and " + "% clauses on % " + "construct"); + } break; case OMP_CLAUSE_ALIGNED: @@ -7902,6 +7934,26 @@ gimplify_omp_for (tree *expr_p, gimple_s OMP_CLAUSE_LINEAR_NO_COPYOUT (c) = 1; flags |= GOVD_LINEAR_LASTPRIVATE_NO_OUTER; } + else + { + struct gimplify_omp_ctx *octx = outer->outer_context; + if (octx + && octx->region_type == ORT_COMBINED_PARALLEL + && octx->outer_context + && (octx->outer_context->region_type + == ORT_WORKSHARE) + && octx->outer_context->combined_loop) + { + octx = octx->outer_context; + n = splay_tree_lookup (octx->variables, + (splay_tree_key)decl); + if (n != NULL && (n->value & GOVD_LOCAL) != 0) + { + OMP_CLAUSE_LINEAR_NO_COPYOUT (c) = 1; + flags |= GOVD_LINEAR_LASTPRIVATE_NO_OUTER; + } + } + } } } @@ -7936,7 +7988,41 @@ gimplify_omp_for (tree *expr_p, gimple_s { omp_add_variable (outer, decl, GOVD_LASTPRIVATE | GOVD_SEEN); - if (outer->outer_context) + if (outer->region_type == ORT_COMBINED_PARALLEL + && outer->outer_context + && (outer->outer_context->region_type + == ORT_WORKSHARE) + && outer->outer_context->combined_loop) + { + outer = outer->outer_context; + n = splay_tree_lookup (outer->variables, + (splay_tree_key)decl); + if (omp_check_private (outer, decl, false)) + outer = NULL; + else if (n == NULL + || ((n->value & GOVD_DATA_SHARE_CLASS) + == 0)) + omp_add_variable (outer, decl, + GOVD_LASTPRIVATE + | GOVD_SEEN); + else + outer = NULL; + } + if (outer && outer->outer_context + && (outer->outer_context->region_type + == ORT_COMBINED_TEAMS)) + { + outer = outer->outer_context; + n = splay_tree_lookup (outer->variables, + (splay_tree_key)decl); + if (n == NULL + || (n->value & GOVD_DATA_SHARE_CLASS) == 0) + omp_add_variable (outer, decl, + GOVD_SHARED | GOVD_SEEN); + else + outer = NULL; + } + if (outer && outer->outer_context) omp_notice_variable (outer->outer_context, decl, true); } @@ -7985,7 +8071,41 @@ gimplify_omp_for (tree *expr_p, gimple_s { omp_add_variable (outer, decl, GOVD_LASTPRIVATE | GOVD_SEEN); - if (outer->outer_context) + if (outer->region_type == ORT_COMBINED_PARALLEL + && outer->outer_context + && (outer->outer_context->region_type + == ORT_WORKSHARE) + && outer->outer_context->combined_loop) + { + outer = outer->outer_context; + n = splay_tree_lookup (outer->variables, + (splay_tree_key)decl); + if (omp_check_private (outer, decl, false)) + outer = NULL; + else if (n == NULL + || ((n->value & GOVD_DATA_SHARE_CLASS) + == 0)) + omp_add_variable (outer, decl, + GOVD_LASTPRIVATE + | GOVD_SEEN); + else + outer = NULL; + } + if (outer && outer->outer_context + && (outer->outer_context->region_type + == ORT_COMBINED_TEAMS)) + { + outer = outer->outer_context; + n = splay_tree_lookup (outer->variables, + (splay_tree_key)decl); + if (n == NULL + || (n->value & GOVD_DATA_SHARE_CLASS) == 0) + omp_add_variable (outer, decl, + GOVD_SHARED | GOVD_SEEN); + else + outer = NULL; + } + if (outer && outer->outer_context) omp_notice_variable (outer->outer_context, decl, true); } --- gcc/omp-low.c.jj 2015-10-22 14:11:23.488003543 +0200 +++ gcc/omp-low.c 2015-10-23 13:10:16.995777473 +0200 @@ -2646,12 +2646,15 @@ add_taskreg_looptemp_clauses (enum gf_ma && TREE_CODE (fd.loop.n2) != INTEGER_CST) { count += fd.collapse - 1; - /* For taskloop, if there are lastprivate clauses on the inner + /* If there are lastprivate clauses on the inner GIMPLE_OMP_FOR, add one more temporaries for the total number of iterations (product of count1 ... countN-1). */ - if (msk == GF_OMP_FOR_KIND_TASKLOOP - && find_omp_clause (gimple_omp_for_clauses (for_stmt), - OMP_CLAUSE_LASTPRIVATE)) + if (find_omp_clause (gimple_omp_for_clauses (for_stmt), + OMP_CLAUSE_LASTPRIVATE)) + count++; + else if (msk == GF_OMP_FOR_KIND_FOR + && find_omp_clause (gimple_omp_parallel_clauses (stmt), + OMP_CLAUSE_LASTPRIVATE)) count++; } for (i = 0; i < count; i++) @@ -8648,6 +8651,30 @@ expand_omp_for_static_nochunk (struct om OMP_CLAUSE__LOOPTEMP_); gcc_assert (innerc); endvar = OMP_CLAUSE_DECL (innerc); + if (fd->collapse > 1 && TREE_CODE (fd->loop.n2) != INTEGER_CST + && gimple_omp_for_kind (fd->for_stmt) == GF_OMP_FOR_KIND_DISTRIBUTE) + { + int i; + for (i = 1; i < fd->collapse; i++) + { + innerc = find_omp_clause (OMP_CLAUSE_CHAIN (innerc), + OMP_CLAUSE__LOOPTEMP_); + gcc_assert (innerc); + } + innerc = find_omp_clause (OMP_CLAUSE_CHAIN (innerc), + OMP_CLAUSE__LOOPTEMP_); + if (innerc) + { + /* If needed (distribute parallel for with lastprivate), + propagate down the total number of iterations. */ + tree t = fold_convert (TREE_TYPE (OMP_CLAUSE_DECL (innerc)), + fd->loop.n2); + t = force_gimple_operand_gsi (&gsi, t, false, NULL_TREE, false, + GSI_CONTINUE_LINKING); + assign_stmt = gimple_build_assign (OMP_CLAUSE_DECL (innerc), t); + gsi_insert_after (&gsi, assign_stmt, GSI_CONTINUE_LINKING); + } + } } t = fold_convert (itype, s0); t = fold_build2 (MULT_EXPR, itype, t, step); @@ -9133,6 +9160,30 @@ expand_omp_for_static_chunk (struct omp_ OMP_CLAUSE__LOOPTEMP_); gcc_assert (innerc); endvar = OMP_CLAUSE_DECL (innerc); + if (fd->collapse > 1 && TREE_CODE (fd->loop.n2) != INTEGER_CST + && gimple_omp_for_kind (fd->for_stmt) == GF_OMP_FOR_KIND_DISTRIBUTE) + { + int i; + for (i = 1; i < fd->collapse; i++) + { + innerc = find_omp_clause (OMP_CLAUSE_CHAIN (innerc), + OMP_CLAUSE__LOOPTEMP_); + gcc_assert (innerc); + } + innerc = find_omp_clause (OMP_CLAUSE_CHAIN (innerc), + OMP_CLAUSE__LOOPTEMP_); + if (innerc) + { + /* If needed (distribute parallel for with lastprivate), + propagate down the total number of iterations. */ + tree t = fold_convert (TREE_TYPE (OMP_CLAUSE_DECL (innerc)), + fd->loop.n2); + t = force_gimple_operand_gsi (&gsi, t, false, NULL_TREE, false, + GSI_CONTINUE_LINKING); + assign_stmt = gimple_build_assign (OMP_CLAUSE_DECL (innerc), t); + gsi_insert_after (&gsi, assign_stmt, GSI_CONTINUE_LINKING); + } + } } t = fold_convert (itype, s0); @@ -13456,26 +13507,36 @@ lower_omp_for_lastprivate (struct omp_fo && TREE_CODE (n2) != INTEGER_CST && gimple_omp_for_combined_into_p (fd->for_stmt)) { - struct omp_context *task_ctx = NULL; + struct omp_context *taskreg_ctx = NULL; if (gimple_code (ctx->outer->stmt) == GIMPLE_OMP_FOR) { gomp_for *gfor = as_a (ctx->outer->stmt); - if (gimple_omp_for_kind (gfor) == GF_OMP_FOR_KIND_FOR) + if (gimple_omp_for_kind (gfor) == GF_OMP_FOR_KIND_FOR + || gimple_omp_for_kind (gfor) == GF_OMP_FOR_KIND_DISTRIBUTE) { - struct omp_for_data outer_fd; - extract_omp_for_data (gfor, &outer_fd, NULL); - n2 = fold_convert (TREE_TYPE (n2), outer_fd.loop.n2); + if (gimple_omp_for_combined_into_p (gfor)) + { + gcc_assert (ctx->outer->outer + && is_parallel_ctx (ctx->outer->outer)); + taskreg_ctx = ctx->outer->outer; + } + else + { + struct omp_for_data outer_fd; + extract_omp_for_data (gfor, &outer_fd, NULL); + n2 = fold_convert (TREE_TYPE (n2), outer_fd.loop.n2); + } } else if (gimple_omp_for_kind (gfor) == GF_OMP_FOR_KIND_TASKLOOP) - task_ctx = ctx->outer->outer; + taskreg_ctx = ctx->outer->outer; } - else if (is_task_ctx (ctx->outer)) - task_ctx = ctx->outer; - if (task_ctx) + else if (is_taskreg_ctx (ctx->outer)) + taskreg_ctx = ctx->outer; + if (taskreg_ctx) { int i; tree innerc - = find_omp_clause (gimple_omp_task_clauses (task_ctx->stmt), + = find_omp_clause (gimple_omp_taskreg_clauses (taskreg_ctx->stmt), OMP_CLAUSE__LOOPTEMP_); gcc_assert (innerc); for (i = 0; i < fd->collapse; i++) @@ -13489,7 +13550,7 @@ lower_omp_for_lastprivate (struct omp_fo if (innerc) n2 = fold_convert (TREE_TYPE (n2), lookup_decl (OMP_CLAUSE_DECL (innerc), - task_ctx)); + taskreg_ctx)); } } cond = build2 (cond_code, boolean_type_node, fd->loop.v, n2); --- gcc/c-family/c-omp.c.jj 2015-10-22 13:41:23.256881650 +0200 +++ gcc/c-family/c-omp.c 2015-10-22 14:35:41.075032703 +0200 @@ -1097,10 +1097,24 @@ c_omp_split_clauses (location_t loc, enu s = C_OMP_CLAUSE_SPLIT_FOR; } break; - /* Lastprivate is allowed on for, sections and simd. In + /* Lastprivate is allowed on distribute, for, sections and simd. In parallel {for{, simd},sections} we actually want to put it on parallel rather than for or sections. */ case OMP_CLAUSE_LASTPRIVATE: + if (code == OMP_DISTRIBUTE) + { + s = C_OMP_CLAUSE_SPLIT_DISTRIBUTE; + break; + } + if ((mask & (OMP_CLAUSE_MASK_1 + << PRAGMA_OMP_CLAUSE_DIST_SCHEDULE)) != 0) + { + c = build_omp_clause (OMP_CLAUSE_LOCATION (clauses), + OMP_CLAUSE_LASTPRIVATE); + OMP_CLAUSE_DECL (c) = OMP_CLAUSE_DECL (clauses); + OMP_CLAUSE_CHAIN (c) = cclauses[C_OMP_CLAUSE_SPLIT_DISTRIBUTE]; + cclauses[C_OMP_CLAUSE_SPLIT_DISTRIBUTE] = c; + } if (code == OMP_FOR || code == OMP_SECTIONS) { if ((mask & (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NUM_THREADS)) --- gcc/c/c-parser.c.jj 2015-10-22 13:41:23.230882028 +0200 +++ gcc/c/c-parser.c 2015-10-22 14:35:41.075032703 +0200 @@ -14682,6 +14682,7 @@ c_parser_omp_cancellation_point (c_parse #define OMP_DISTRIBUTE_CLAUSE_MASK \ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRIVATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_FIRSTPRIVATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_LASTPRIVATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DIST_SCHEDULE)\ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COLLAPSE)) --- gcc/cp/parser.c.jj 2015-10-22 13:41:23.248881766 +0200 +++ gcc/cp/parser.c 2015-10-22 14:35:41.067032818 +0200 @@ -33572,6 +33572,7 @@ cp_parser_omp_cancellation_point (cp_par #define OMP_DISTRIBUTE_CLAUSE_MASK \ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRIVATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_FIRSTPRIVATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_LASTPRIVATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DIST_SCHEDULE)\ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COLLAPSE)) --- gcc/testsuite/c-c++-common/gomp/distribute-1.c.jj 2015-10-22 14:35:41.068032804 +0200 +++ gcc/testsuite/c-c++-common/gomp/distribute-1.c 2015-10-22 14:35:41.068032804 +0200 @@ -0,0 +1,56 @@ +int s1, s2, s3, s4, s5, s6, s7, s8; +#pragma omp declare target (s1, s2, s3, s4, s5, s6, s7, s8) + +void +f1 (void) +{ + int i; + #pragma omp distribute + for (i = 0; i < 64; i++) + ; + #pragma omp distribute private (i) + for (i = 0; i < 64; i++) + ; + #pragma omp distribute + for (int j = 0; j < 64; j++) + ; + #pragma omp distribute lastprivate (s1) + for (s1 = 0; s1 < 64; s1 += 2) + ; + #pragma omp distribute lastprivate (s2) + for (i = 0; i < 64; i++) + s2 = 2 * i; + #pragma omp distribute simd + for (i = 0; i < 64; i++) + ; + #pragma omp distribute simd lastprivate (s3, s4) collapse(2) + for (s3 = 0; s3 < 64; s3++) + for (s4 = 0; s4 < 3; s4++) + ; + #pragma omp distribute parallel for + for (i = 0; i < 64; i++) + ; + #pragma omp distribute parallel for private (i) + for (i = 0; i < 64; i++) + ; + #pragma omp distribute parallel for lastprivate (s5) + for (s5 = 0; s5 < 64; s5++) + ; + #pragma omp distribute firstprivate (s7) private (s8) + for (i = 0; i < 64; i++) + s8 = s7++; +} + +void +f2 (void) +{ + int i; + #pragma omp distribute lastprivate (i) /* { dg-error "lastprivate variable .i. is private in outer context" } */ + for (i = 0; i < 64; i++) + ; + #pragma omp distribute firstprivate (s6) lastprivate (s6) /* { dg-error "same variable used in .firstprivate. and .lastprivate. clauses on .distribute. construct" } */ + for (i = 0; i < 64; i++) + s6 += i; +} + +#pragma omp declare target to(f1, f2) --- gcc/testsuite/c-c++-common/gomp/pr61486-2.c.jj 2015-10-22 13:41:22.983885625 +0200 +++ gcc/testsuite/c-c++-common/gomp/pr61486-2.c 2015-10-22 14:35:41.068032804 +0200 @@ -355,8 +355,11 @@ test (int n, int o, int p, int q, int r, int q, i, j; +#pragma omp declare target +int s; + void -test2 (int n, int o, int p, int r, int s, int *pp) +test2 (int n, int o, int p, int r, int *pp) { int a[o]; #pragma omp distribute collapse (2) dist_schedule (static, 4) firstprivate (q) @@ -449,3 +452,4 @@ test2 (int n, int o, int p, int r, int s s = i * 10; } } +#pragma omp end declare target --- libgomp/testsuite/libgomp.c/pr66199-5.c.jj 2015-10-22 14:35:41.080032631 +0200 +++ libgomp/testsuite/libgomp.c/pr66199-5.c 2015-10-22 16:15:11.064658921 +0200 @@ -0,0 +1,66 @@ +/* PR middle-end/66199 */ +/* { dg-do run } */ + +#pragma omp declare target +int u[1024], v[1024], w[1024]; +#pragma omp end declare target + +__attribute__((noinline, noclone)) long +f1 (long a, long b) +{ + long d; + #pragma omp target map(from: d) + #pragma omp teams distribute parallel for simd default(none) firstprivate (a, b) shared(u, v, w) + for (d = a; d < b; d++) + u[d] = v[d] + w[d]; + return d; +} + +__attribute__((noinline, noclone)) long +f2 (long a, long b, long c) +{ + long d, e; + #pragma omp target map(from: d, e) + #pragma omp teams distribute parallel for simd default(none) firstprivate (a, b, c) shared(u, v, w) linear(d) lastprivate(e) + for (d = a; d < b; d++) + { + u[d] = v[d] + w[d]; + e = c + d * 5; + } + return d + e; +} + +__attribute__((noinline, noclone)) long +f3 (long a1, long b1, long a2, long b2) +{ + long d1, d2; + #pragma omp target map(from: d1, d2) + #pragma omp teams distribute parallel for simd default(none) firstprivate (a1, b1, a2, b2) shared(u, v, w) lastprivate(d1, d2) collapse(2) + for (d1 = a1; d1 < b1; d1++) + for (d2 = a2; d2 < b2; d2++) + u[d1 * 32 + d2] = v[d1 * 32 + d2] + w[d1 * 32 + d2]; + return d1 + d2; +} + +__attribute__((noinline, noclone)) long +f4 (long a1, long b1, long a2, long b2) +{ + long d1, d2; + #pragma omp target map(from: d1, d2) + #pragma omp teams distribute parallel for simd default(none) firstprivate (a1, b1, a2, b2) shared(u, v, w) collapse(2) + for (d1 = a1; d1 < b1; d1++) + for (d2 = a2; d2 < b2; d2++) + u[d1 * 32 + d2] = v[d1 * 32 + d2] + w[d1 * 32 + d2]; + return d1 + d2; +} + +int +main () +{ + if (f1 (0, 1024) != 1024 + || f2 (0, 1024, 17) != 1024 + (17 + 5 * 1023) + || f3 (0, 32, 0, 32) != 64 + || f4 (0, 32, 0, 32) != 64) + __builtin_abort (); + return 0; +} --- libgomp/testsuite/libgomp.c/pr66199-6.c.jj 2015-10-22 14:35:41.080032631 +0200 +++ libgomp/testsuite/libgomp.c/pr66199-6.c 2015-10-22 16:24:37.000000000 +0200 @@ -0,0 +1,42 @@ +/* PR middle-end/66199 */ +/* { dg-do run } */ +/* { dg-options "-O2 -fopenmp" } */ + +#pragma omp declare target +int u[1024], v[1024], w[1024]; +#pragma omp end declare target + +__attribute__((noinline, noclone)) long +f2 (long a, long b, long c) +{ + long d, e; + #pragma omp target map(from: d, e) + #pragma omp teams distribute parallel for default(none) firstprivate (a, b, c) shared(u, v, w) lastprivate(d, e) + for (d = a; d < b; d++) + { + u[d] = v[d] + w[d]; + e = c + d * 5; + } + return d + e; +} + +__attribute__((noinline, noclone)) long +f3 (long a1, long b1, long a2, long b2) +{ + long d1, d2; + #pragma omp target map(from: d1, d2) + #pragma omp teams distribute parallel for default(none) firstprivate (a1, b1, a2, b2) shared(u, v, w) lastprivate(d1, d2) collapse(2) + for (d1 = a1; d1 < b1; d1++) + for (d2 = a2; d2 < b2; d2++) + u[d1 * 32 + d2] = v[d1 * 32 + d2] + w[d1 * 32 + d2]; + return d1 + d2; +} + +int +main () +{ + if (f2 (0, 1024, 17) != 1024 + (17 + 5 * 1023) + || f3 (0, 32, 0, 32) != 64) + __builtin_abort (); + return 0; +} --- libgomp/testsuite/libgomp.c/pr66199-7.c.jj 2015-10-22 14:35:41.080032631 +0200 +++ libgomp/testsuite/libgomp.c/pr66199-7.c 2015-10-22 17:11:33.419476232 +0200 @@ -0,0 +1,66 @@ +/* PR middle-end/66199 */ +/* { dg-do run } */ + +#pragma omp declare target +int u[1024], v[1024], w[1024]; +#pragma omp end declare target + +__attribute__((noinline, noclone)) long +f1 (long a, long b) +{ + long d; + #pragma omp target map(from: d) + #pragma omp teams distribute simd default(none) firstprivate (a, b) shared(u, v, w) + for (d = a; d < b; d++) + u[d] = v[d] + w[d]; + return d; +} + +__attribute__((noinline, noclone)) long +f2 (long a, long b, long c) +{ + long d, e; + #pragma omp target map(from: d, e) + #pragma omp teams distribute simd default(none) firstprivate (a, b, c) shared(u, v, w) linear(d) lastprivate(e) + for (d = a; d < b; d++) + { + u[d] = v[d] + w[d]; + e = c + d * 5; + } + return d + e; +} + +__attribute__((noinline, noclone)) long +f3 (long a1, long b1, long a2, long b2) +{ + long d1, d2; + #pragma omp target map(from: d1, d2) + #pragma omp teams distribute simd default(none) firstprivate (a1, b1, a2, b2) shared(u, v, w) lastprivate(d1, d2) collapse(2) + for (d1 = a1; d1 < b1; d1++) + for (d2 = a2; d2 < b2; d2++) + u[d1 * 32 + d2] = v[d1 * 32 + d2] + w[d1 * 32 + d2]; + return d1 + d2; +} + +__attribute__((noinline, noclone)) long +f4 (long a1, long b1, long a2, long b2) +{ + long d1, d2; + #pragma omp target map(from: d1, d2) + #pragma omp teams distribute simd default(none) firstprivate (a1, b1, a2, b2) shared(u, v, w) collapse(2) + for (d1 = a1; d1 < b1; d1++) + for (d2 = a2; d2 < b2; d2++) + u[d1 * 32 + d2] = v[d1 * 32 + d2] + w[d1 * 32 + d2]; + return d1 + d2; +} + +int +main () +{ + if (f1 (0, 1024) != 1024 + || f2 (0, 1024, 17) != 1024 + (17 + 5 * 1023) + || f3 (0, 32, 0, 32) != 64 + || f4 (0, 32, 0, 32) != 64) + __builtin_abort (); + return 0; +} --- libgomp/testsuite/libgomp.c/pr66199-8.c.jj 2015-10-22 14:35:41.080032631 +0200 +++ libgomp/testsuite/libgomp.c/pr66199-8.c 2015-10-22 17:11:33.419476232 +0200 @@ -0,0 +1,70 @@ +/* PR middle-end/66199 */ +/* { dg-do run } */ + +#pragma omp declare target +int u[1024], v[1024], w[1024]; +#pragma omp end declare target + +__attribute__((noinline, noclone)) long +f1 (long a, long b) +{ + long d; + #pragma omp target map(from: d) + #pragma omp teams default(none) shared(a, b, d, u, v, w) + #pragma omp distribute simd firstprivate (a, b) + for (d = a; d < b; d++) + u[d] = v[d] + w[d]; + return d; +} + +__attribute__((noinline, noclone)) long +f2 (long a, long b, long c) +{ + long d, e; + #pragma omp target map(from: d, e) + #pragma omp teams default(none) firstprivate (a, b, c) shared(d, e, u, v, w) + #pragma omp distribute simd linear(d) lastprivate(e) + for (d = a; d < b; d++) + { + u[d] = v[d] + w[d]; + e = c + d * 5; + } + return d + e; +} + +__attribute__((noinline, noclone)) long +f3 (long a1, long b1, long a2, long b2) +{ + long d1, d2; + #pragma omp target map(from: d1, d2) + #pragma omp teams default(none) shared(a1, b1, a2, b2, d1, d2, u, v, w) + #pragma omp distribute simd firstprivate (a1, b1, a2, b2) lastprivate(d1, d2) collapse(2) + for (d1 = a1; d1 < b1; d1++) + for (d2 = a2; d2 < b2; d2++) + u[d1 * 32 + d2] = v[d1 * 32 + d2] + w[d1 * 32 + d2]; + return d1 + d2; +} + +__attribute__((noinline, noclone)) long +f4 (long a1, long b1, long a2, long b2) +{ + long d1, d2; + #pragma omp target map(from: d1, d2) + #pragma omp teams default(none) firstprivate (a1, b1, a2, b2) shared(d1, d2, u, v, w) + #pragma omp distribute simd collapse(2) + for (d1 = a1; d1 < b1; d1++) + for (d2 = a2; d2 < b2; d2++) + u[d1 * 32 + d2] = v[d1 * 32 + d2] + w[d1 * 32 + d2]; + return d1 + d2; +} + +int +main () +{ + if (f1 (0, 1024) != 1024 + || f2 (0, 1024, 17) != 1024 + (17 + 5 * 1023) + || f3 (0, 32, 0, 32) != 64 + || f4 (0, 32, 0, 32) != 64) + __builtin_abort (); + return 0; +} --- libgomp/testsuite/libgomp.c/pr66199-9.c.jj 2015-10-22 17:11:56.362148829 +0200 +++ libgomp/testsuite/libgomp.c/pr66199-9.c 2015-10-22 17:12:59.597246434 +0200 @@ -0,0 +1,43 @@ +/* PR middle-end/66199 */ +/* { dg-do run } */ + +#pragma omp declare target +int u[1024], v[1024], w[1024]; +#pragma omp end declare target + +__attribute__((noinline, noclone)) long +f2 (long a, long b, long c) +{ + long d, e; + #pragma omp target map(from: d, e) + #pragma omp teams default(none) firstprivate (a, b, c) shared(d, e, u, v, w) + #pragma omp distribute lastprivate(d, e) + for (d = a; d < b; d++) + { + u[d] = v[d] + w[d]; + e = c + d * 5; + } + return d + e; +} + +__attribute__((noinline, noclone)) long +f3 (long a1, long b1, long a2, long b2) +{ + long d1, d2; + #pragma omp target map(from: d1, d2) + #pragma omp teams default(none) shared(a1, b1, a2, b2, d1, d2, u, v, w) + #pragma omp distribute firstprivate (a1, b1, a2, b2) lastprivate(d1, d2) collapse(2) + for (d1 = a1; d1 < b1; d1++) + for (d2 = a2; d2 < b2; d2++) + u[d1 * 32 + d2] = v[d1 * 32 + d2] + w[d1 * 32 + d2]; + return d1 + d2; +} + +int +main () +{ + if (f2 (0, 1024, 17) != 1024 + (17 + 5 * 1023) + || f3 (0, 32, 0, 32) != 64) + __builtin_abort (); + return 0; +} --- libgomp/testsuite/libgomp.c++/pr66199-3.C.jj 2015-10-22 17:14:28.277980917 +0200 +++ libgomp/testsuite/libgomp.c++/pr66199-3.C 2015-10-22 17:14:28.276980932 +0200 @@ -0,0 +1,4 @@ +// PR middle-end/66199 +// { dg-do run } + +#include "../libgomp.c/pr66199-3.c" --- libgomp/testsuite/libgomp.c++/pr66199-4.C.jj 2015-10-22 17:14:28.280980875 +0200 +++ libgomp/testsuite/libgomp.c++/pr66199-4.C 2015-10-22 17:14:28.279980889 +0200 @@ -0,0 +1,4 @@ +// PR middle-end/66199 +// { dg-do run } + +#include "../libgomp.c/pr66199-4.c" --- libgomp/testsuite/libgomp.c++/pr66199-5.C.jj 2015-10-22 17:14:28.283980832 +0200 +++ libgomp/testsuite/libgomp.c++/pr66199-5.C 2015-10-22 17:14:28.282980846 +0200 @@ -0,0 +1,4 @@ +// PR middle-end/66199 +// { dg-do run } + +#include "../libgomp.c/pr66199-5.c" --- libgomp/testsuite/libgomp.c++/pr66199-6.C.jj 2015-10-22 17:14:28.287980775 +0200 +++ libgomp/testsuite/libgomp.c++/pr66199-6.C 2015-10-22 17:14:28.285980803 +0200 @@ -0,0 +1,4 @@ +// PR middle-end/66199 +// { dg-do run } + +#include "../libgomp.c/pr66199-6.c" --- libgomp/testsuite/libgomp.c++/pr66199-7.C.jj 2015-10-22 17:14:28.290980732 +0200 +++ libgomp/testsuite/libgomp.c++/pr66199-7.C 2015-10-22 17:14:28.289980746 +0200 @@ -0,0 +1,4 @@ +// PR middle-end/66199 +// { dg-do run } + +#include "../libgomp.c/pr66199-7.c" --- libgomp/testsuite/libgomp.c++/pr66199-8.C.jj 2015-10-22 17:14:28.293980689 +0200 +++ libgomp/testsuite/libgomp.c++/pr66199-8.C 2015-10-22 17:14:28.292980703 +0200 @@ -0,0 +1,4 @@ +// PR middle-end/66199 +// { dg-do run } + +#include "../libgomp.c/pr66199-8.c" --- libgomp/testsuite/libgomp.c++/pr66199-9.C.jj 2015-10-22 17:14:28.296980646 +0200 +++ libgomp/testsuite/libgomp.c++/pr66199-9.C 2015-10-22 17:14:28.295980661 +0200 @@ -0,0 +1,4 @@ +// PR middle-end/66199 +// { dg-do run } + +#include "../libgomp.c/pr66199-9.c"