From patchwork Thu Mar 20 14:42:48 2014 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 332198 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 547972C009A for ; Fri, 21 Mar 2014 01:43:25 +1100 (EST) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:in-reply-to:references:date:message-id :mime-version:content-type; q=dns; s=default; b=xou7ozdcLPXOkSM2 HivUM6dZ0jrJ3rR+9l5EARmMp51oV0BtA7LOetLOx6vYkSg4bqsvqVMNsQjQOs6v X8edztpWZbEGiHmmuuVObq7aJCquVzYwpNp5JL27y+YnSUI5Y2NgzUyDAdjJZ4KV aKiN8K1SpluJGBWXt8+SEytEV7k= 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:from :to:cc:subject:in-reply-to:references:date:message-id :mime-version:content-type; s=default; bh=58jhoz6pcTrylVVH7vSTf+ MGEO0=; b=oiSt92xW28bdqmCaOJrrrNIWC9Rg+PCOsyUfTdER6FsU83UeaHgmPG Msj9sHM3Y5WeXIv1nE3DunAzcPNhOya8PYAq3oOjJpEHOMljzgd5QtBoB59gjsOj 9tkn3KZNCWXpAW11M61DaaUXjEqtG9QLXdfjK93xS1O5WDki1Ocsw= Received: (qmail 1523 invoked by alias); 20 Mar 2014 14:43:15 -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 1512 invoked by uid 89); 20 Mar 2014 14:43:15 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-0.7 required=5.0 tests=AWL, BAYES_50 autolearn=ham version=3.3.2 X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Thu, 20 Mar 2014 14:43:09 +0000 Received: from svr-orw-exc-10.mgc.mentorg.com ([147.34.98.58]) by relay1.mentorg.com with esmtp id 1WQeBT-0005eh-9U from Thomas_Schwinge@mentor.com ; Thu, 20 Mar 2014 07:43:03 -0700 Received: from SVR-IES-FEM-01.mgc.mentorg.com ([137.202.0.104]) by SVR-ORW-EXC-10.mgc.mentorg.com with Microsoft SMTPSVC(6.0.3790.4675); Thu, 20 Mar 2014 07:43:03 -0700 Received: from feldtkeller.schwinge.homeip.net (137.202.0.76) by SVR-IES-FEM-01.mgc.mentorg.com (137.202.0.104) with Microsoft SMTP Server id 14.2.247.3; Thu, 20 Mar 2014 14:42:58 +0000 From: Thomas Schwinge To: , Ilmir Usmanov , CC: Slava Garbuzov , Evgeny Gavrin Subject: Re: [PATCH] [gomp4] Initial support of OpenACC loop directive in C front-end. In-Reply-To: <87ha6vipjf.fsf@schwinge.name> References: <53283E04.6010501@samsung.com> <87ha6vipjf.fsf@schwinge.name> User-Agent: Notmuch/0.9-101-g81dad07 (http://notmuchmail.org) Emacs/23.4.1 (i486-pc-linux-gnu) Date: Thu, 20 Mar 2014 15:42:48 +0100 Message-ID: <87zjklkk2f.fsf@kepler.schwinge.homeip.net> MIME-Version: 1.0 Hi! On Tue, 18 Mar 2014 14:50:44 +0100, I wrote: > On Tue, 18 Mar 2014 16:37:24 +0400, Ilmir Usmanov wrote: > > This patch introduces support of OpenACC loop directive (and combined > > directives) in C front-end up to GENERIC. Currently no clause is allowed. > > Thanks! I had worked on a simpler patch, not yet dealing with combined > clauses. Also, I have some work for the GIMPLE level, namely building on > GIMPLE_OMP_FOR, adding a new GF_OMP_FOR_KIND_OACC_LOOP. I'll post this > soon. Here are the patches, committed in r208702..4 to gomp-4_0-branch. Jakub, are the first two fine for trunk, or shall I wait until stage 1? commit 834daebdafa1cf4f8507fb932b7115ba3ebb02a3 Author: tschwinge Date: Thu Mar 20 14:39:30 2014 +0000 Just enumerate all GF_OMP_FOR_KIND_* and GF_OMP_TARGET_KIND_*. gcc/ * gimple.h (enum gf_mask): Rewrite "<< 0" shift expressions used for GF_OMP_FOR_KIND_MASK, GF_OMP_FOR_KIND_FOR, GF_OMP_FOR_KIND_DISTRIBUTE, GF_OMP_FOR_KIND_SIMD, GF_OMP_FOR_KIND_CILKSIMD, GF_OMP_TARGET_KIND_MASK, GF_OMP_TARGET_KIND_REGION, GF_OMP_TARGET_KIND_DATA, GF_OMP_TARGET_KIND_UPDATE, GF_OMP_TARGET_KIND_OACC_DATA. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@208702 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/ChangeLog.gomp | 7 +++++++ gcc/gimple.h | 20 ++++++++++---------- 2 files changed, 17 insertions(+), 10 deletions(-) Grüße, Thomas diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp index f43452c..72828fb 100644 --- gcc/ChangeLog.gomp +++ gcc/ChangeLog.gomp @@ -1,5 +1,12 @@ 2014-03-20 Thomas Schwinge + * gimple.h (enum gf_mask): Rewrite "<< 0" shift expressions used + for GF_OMP_FOR_KIND_MASK, GF_OMP_FOR_KIND_FOR, + GF_OMP_FOR_KIND_DISTRIBUTE, GF_OMP_FOR_KIND_SIMD, + GF_OMP_FOR_KIND_CILKSIMD, GF_OMP_TARGET_KIND_MASK, + GF_OMP_TARGET_KIND_REGION, GF_OMP_TARGET_KIND_DATA, + GF_OMP_TARGET_KIND_UPDATE, GF_OMP_TARGET_KIND_OACC_DATA. + * omp-low.c (check_omp_nesting_restrictions): Allow nesting of OpenACC constructs inside of OpenACC data constructs. diff --git gcc/gimple.h gcc/gimple.h index 910072d..17441ac 100644 --- gcc/gimple.h +++ gcc/gimple.h @@ -91,18 +91,18 @@ enum gf_mask { GF_CALL_ALLOCA_FOR_VAR = 1 << 5, GF_CALL_INTERNAL = 1 << 6, GF_OMP_PARALLEL_COMBINED = 1 << 0, - GF_OMP_FOR_KIND_MASK = 3 << 0, - GF_OMP_FOR_KIND_FOR = 0 << 0, - GF_OMP_FOR_KIND_DISTRIBUTE = 1 << 0, - GF_OMP_FOR_KIND_SIMD = 2 << 0, - GF_OMP_FOR_KIND_CILKSIMD = 3 << 0, + GF_OMP_FOR_KIND_MASK = (1 << 2) - 1, + GF_OMP_FOR_KIND_FOR = 0, + GF_OMP_FOR_KIND_DISTRIBUTE = 1, + GF_OMP_FOR_KIND_SIMD = 2, + GF_OMP_FOR_KIND_CILKSIMD = 3, GF_OMP_FOR_COMBINED = 1 << 2, GF_OMP_FOR_COMBINED_INTO = 1 << 3, - GF_OMP_TARGET_KIND_MASK = 3 << 0, - GF_OMP_TARGET_KIND_REGION = 0 << 0, - GF_OMP_TARGET_KIND_DATA = 1 << 0, - GF_OMP_TARGET_KIND_UPDATE = 2 << 0, - GF_OMP_TARGET_KIND_OACC_DATA = 3 << 0, + GF_OMP_TARGET_KIND_MASK = (1 << 2) - 1, + GF_OMP_TARGET_KIND_REGION = 0, + GF_OMP_TARGET_KIND_DATA = 1, + GF_OMP_TARGET_KIND_UPDATE = 2, + GF_OMP_TARGET_KIND_OACC_DATA = 3, /* True on an GIMPLE_OMP_RETURN statement if the return does not require a thread synchronization via some sort of barrier. The exact barrier commit c32a48d3d47bbaa811991e2e5f42e62d9c715a60 Author: tschwinge Date: Thu Mar 20 14:39:42 2014 +0000 GF_OMP_FOR_SIMD: Flag for SIMD variants of OMP_FOR kinds. gcc/ * gimple.h (enum gf_mask): Add and use GF_OMP_FOR_SIMD. * omp-low.c: Update accordingly. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@208703 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/ChangeLog.gomp | 3 +++ gcc/gimple.h | 6 ++++-- gcc/omp-low.c | 18 +++++++++--------- 3 files changed, 16 insertions(+), 11 deletions(-) diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp index 72828fb..1753d73 100644 --- gcc/ChangeLog.gomp +++ gcc/ChangeLog.gomp @@ -1,5 +1,8 @@ 2014-03-20 Thomas Schwinge + * gimple.h (enum gf_mask): Add and use GF_OMP_FOR_SIMD. + * omp-low.c: Update accordingly. + * gimple.h (enum gf_mask): Rewrite "<< 0" shift expressions used for GF_OMP_FOR_KIND_MASK, GF_OMP_FOR_KIND_FOR, GF_OMP_FOR_KIND_DISTRIBUTE, GF_OMP_FOR_KIND_SIMD, diff --git gcc/gimple.h gcc/gimple.h index 17441ac..34a0bdb 100644 --- gcc/gimple.h +++ gcc/gimple.h @@ -94,8 +94,10 @@ enum gf_mask { GF_OMP_FOR_KIND_MASK = (1 << 2) - 1, GF_OMP_FOR_KIND_FOR = 0, GF_OMP_FOR_KIND_DISTRIBUTE = 1, - GF_OMP_FOR_KIND_SIMD = 2, - GF_OMP_FOR_KIND_CILKSIMD = 3, + /* Flag for SIMD variants of OMP_FOR kinds. */ + GF_OMP_FOR_SIMD = 1 << 1, + GF_OMP_FOR_KIND_SIMD = GF_OMP_FOR_SIMD | 0, + GF_OMP_FOR_KIND_CILKSIMD = GF_OMP_FOR_SIMD | 1, GF_OMP_FOR_COMBINED = 1 << 2, GF_OMP_FOR_COMBINED_INTO = 1 << 3, GF_OMP_TARGET_KIND_MASK = (1 << 2) - 1, diff --git gcc/omp-low.c gcc/omp-low.c index 23a0dda..c3b3e95 100644 --- gcc/omp-low.c +++ gcc/omp-low.c @@ -298,7 +298,7 @@ extract_omp_for_data (gimple for_stmt, struct omp_for_data *fd, int i; struct omp_for_data_loop dummy_loop; location_t loc = gimple_location (for_stmt); - bool simd = gimple_omp_for_kind (for_stmt) & GF_OMP_FOR_KIND_SIMD; + bool simd = gimple_omp_for_kind (for_stmt) & GF_OMP_FOR_SIMD; bool distribute = gimple_omp_for_kind (for_stmt) == GF_OMP_FOR_KIND_DISTRIBUTE; @@ -1024,7 +1024,7 @@ build_outer_var_ref (tree var, omp_context *ctx) x = build_receiver_ref (var, by_ref, ctx); } else if (gimple_code (ctx->stmt) == GIMPLE_OMP_FOR - && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_KIND_SIMD) + && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_SIMD) { /* #pragma omp simd isn't a worksharing construct, and can reference even private vars in its linear etc. clauses. */ @@ -2451,7 +2451,7 @@ check_omp_nesting_restrictions (gimple stmt, omp_context *ctx) if (ctx != NULL) { if (gimple_code (ctx->stmt) == GIMPLE_OMP_FOR - && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_KIND_SIMD) + && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_SIMD) { error_at (gimple_location (stmt), "OpenMP constructs may not be nested inside simd region"); @@ -2474,7 +2474,7 @@ check_omp_nesting_restrictions (gimple stmt, omp_context *ctx) switch (gimple_code (stmt)) { case GIMPLE_OMP_FOR: - if (gimple_omp_for_kind (stmt) & GF_OMP_FOR_KIND_SIMD) + if (gimple_omp_for_kind (stmt) & GF_OMP_FOR_SIMD) return true; if (gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_DISTRIBUTE) { @@ -2802,7 +2802,7 @@ scan_omp_1_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p, if (setjmp_or_longjmp_p (fndecl) && ctx && gimple_code (ctx->stmt) == GIMPLE_OMP_FOR - && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_KIND_SIMD) + && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_SIMD) { remove = true; error_at (gimple_location (stmt), @@ -3225,7 +3225,7 @@ lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist, bool reduction_omp_orig_ref = false; int pass; bool is_simd = (gimple_code (ctx->stmt) == GIMPLE_OMP_FOR - && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_KIND_SIMD); + && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_SIMD); int max_vf = 0; tree lane = NULL_TREE, idx = NULL_TREE; tree ivar = NULL_TREE, lvar = NULL_TREE; @@ -3969,7 +3969,7 @@ lower_lastprivate_clauses (tree clauses, tree predicate, gimple_seq *stmt_list, } if (gimple_code (ctx->stmt) == GIMPLE_OMP_FOR - && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_KIND_SIMD) + && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_SIMD) { simduid = find_omp_clause (orig_clauses, OMP_CLAUSE__SIMDUID_); if (simduid) @@ -4066,7 +4066,7 @@ lower_reduction_clauses (tree clauses, gimple_seq *stmt_seqp, omp_context *ctx) /* SIMD reductions are handled in lower_rec_input_clauses. */ if (gimple_code (ctx->stmt) == GIMPLE_OMP_FOR - && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_KIND_SIMD) + && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_SIMD) return; /* First see if there is exactly one reduction clause. Use OMP_ATOMIC @@ -7393,7 +7393,7 @@ expand_omp_for (struct omp_region *region, gimple inner_stmt) original loops from being detected. Fix that up. */ loops_state_set (LOOPS_NEED_FIXUP); - if (gimple_omp_for_kind (fd.for_stmt) & GF_OMP_FOR_KIND_SIMD) + if (gimple_omp_for_kind (fd.for_stmt) & GF_OMP_FOR_SIMD) expand_omp_simd (region, &fd); else if (fd.sched_kind == OMP_CLAUSE_SCHEDULE_STATIC && !fd.have_ordered) commit f1d39706db8dccbc988e2c66552511cd54632257 Author: tschwinge Date: Thu Mar 20 14:40:01 2014 +0000 Continue implementation of OpenACC loop construct. gcc/ * gimple.h (enum gf_mask): Add GF_OMP_FOR_KIND_OACC_LOOP. (is_gimple_omp_oacc_specifically): Handle it. * gimple-pretty-print.c (dump_gimple_omp_for): Likewise. * gimple.def (GIMPLE_OMP_FOR): Update for OpenACC loop. * gimple.c (gimple_build_omp_for): Don't explicitly mention some clauses. (gimple_copy) : Handle GF_OMP_FOR_KIND_OACC_LOOP. * omp-low.c (extract_omp_for_data, scan_sharing_clauses) (check_omp_nesting_restrictions, lower_rec_input_clauses) (lower_lastprivate_clauses, lower_reduction_clauses) (expand_omp_for_generic, expand_omp_for_static_nochunk) (expand_omp_for_static_chunk, maybe_add_implicit_barrier_cancel) (lower_omp_for): Likewise. * tree-inline.c (remap_gimple_stmt): Likewise. * tree-nested.c (walk_gimple_omp_for) (convert_nonlocal_reference_stmt, convert_local_reference_stmt) (convert_gimple_call): Likewise. * doc/gimple.texi (GIMPLE_OMP_FOR): Don't explicitly mention some clauses. * gimplify.c (gimplify_omp_for, gimplify_expr): Handle OACC_LOOP. gcc/testsuite/ * c-c++-common/goacc-gomp/nesting-1.c: New file. * c-c++-common/goacc-gomp/nesting-fail-1.c: Extend. * c-c++-common/goacc/clauses-fail.c: Likewise. * c-c++-common/goacc/nesting-1.c: Likewise. * gcc.dg/goacc/sb-1.c: Likewise. * gcc.dg/goacc/sb-3.c: New file. gcc/c-family/ * c-omp.c (check_omp_for_incr_expr, c_finish_omp_for): Update comments. * c-pragma.c (oacc_pragmas): Sort PRAGMA_OACC_LOOP alphabetically. * c-pragma.h (enum pragma_kind): Likewise. gcc/c/ * c-parser.c: Update comments. (c_parser_oacc_loop): Move. (c_parser_omp_for_loop): Catch some unsupported cases. (c_parser_omp_construct) : Sort alphabetically. gcc/ * tree.def (OACC_LOOP): Sort after OMP_DISTRIBUTE. * tree.h (OMP_LOOP_CHECK): Update accordingly. * gimplify.c (is_gimple_stmt): Sort OACC_LOOP after OMP_DISTRIBUTE. * tree-pretty-print.c (dump_generic_node): Likewise. * doc/generic.texi (OACC_LOOP): Sort after OACC_HOST_DATA. (OMP_FOR): Fix and extend for OACC_LOOP. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@208704 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/ChangeLog.gomp | 29 ++++ gcc/c-family/ChangeLog.gomp | 7 + gcc/c-family/c-omp.c | 4 +- gcc/c-family/c-pragma.c | 2 +- gcc/c-family/c-pragma.h | 2 +- gcc/c/ChangeLog.gomp | 8 ++ gcc/c/c-parser.c | 80 ++++++----- gcc/doc/generic.texi | 21 +-- gcc/doc/gimple.texi | 5 +- gcc/gimple-pretty-print.c | 18 ++- gcc/gimple.c | 4 +- gcc/gimple.def | 5 +- gcc/gimple.h | 117 ++++++++------- gcc/gimplify.c | 45 ++++-- gcc/omp-low.c | 157 +++++++++++++++------ gcc/testsuite/ChangeLog.gomp | 7 + gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c | 12 ++ .../c-c++-common/goacc-gomp/nesting-fail-1.c | 98 +++++++++++++ gcc/testsuite/c-c++-common/goacc/clauses-fail.c | 6 + gcc/testsuite/c-c++-common/goacc/nesting-1.c | 43 ++++++ gcc/testsuite/gcc.dg/goacc/sb-1.c | 21 +++ gcc/testsuite/gcc.dg/goacc/sb-3.c | 18 +++ gcc/tree-inline.c | 1 + gcc/tree-nested.c | 5 + gcc/tree-pretty-print.c | 8 +- gcc/tree.def | 10 +- gcc/tree.h | 2 +- 27 files changed, 560 insertions(+), 175 deletions(-) diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp index 1753d73..1d35b58 100644 --- gcc/ChangeLog.gomp +++ gcc/ChangeLog.gomp @@ -1,5 +1,34 @@ 2014-03-20 Thomas Schwinge + * gimple.h (enum gf_mask): Add GF_OMP_FOR_KIND_OACC_LOOP. + (is_gimple_omp_oacc_specifically): Handle it. + * gimple-pretty-print.c (dump_gimple_omp_for): Likewise. + * gimple.def (GIMPLE_OMP_FOR): Update for OpenACC loop. + * gimple.c (gimple_build_omp_for): Don't explicitly mention some + clauses. + (gimple_copy) : Handle GF_OMP_FOR_KIND_OACC_LOOP. + * omp-low.c (extract_omp_for_data, scan_sharing_clauses) + (check_omp_nesting_restrictions, lower_rec_input_clauses) + (lower_lastprivate_clauses, lower_reduction_clauses) + (expand_omp_for_generic, expand_omp_for_static_nochunk) + (expand_omp_for_static_chunk, maybe_add_implicit_barrier_cancel) + (lower_omp_for): Likewise. + * tree-inline.c (remap_gimple_stmt): Likewise. + * tree-nested.c (walk_gimple_omp_for) + (convert_nonlocal_reference_stmt, convert_local_reference_stmt) + (convert_gimple_call): Likewise. + * doc/gimple.texi (GIMPLE_OMP_FOR): Don't explicitly mention some + clauses. + * gimplify.c (gimplify_omp_for, gimplify_expr): Handle OACC_LOOP. + + * tree.def (OACC_LOOP): Sort after OMP_DISTRIBUTE. + * tree.h (OMP_LOOP_CHECK): Update accordingly. + * gimplify.c (is_gimple_stmt): Sort OACC_LOOP after + OMP_DISTRIBUTE. + * tree-pretty-print.c (dump_generic_node): Likewise. + * doc/generic.texi (OACC_LOOP): Sort after OACC_HOST_DATA. + (OMP_FOR): Fix and extend for OACC_LOOP. + * gimple.h (enum gf_mask): Add and use GF_OMP_FOR_SIMD. * omp-low.c: Update accordingly. diff --git gcc/c-family/ChangeLog.gomp gcc/c-family/ChangeLog.gomp index b33b365..37ebfe9 100644 --- gcc/c-family/ChangeLog.gomp +++ gcc/c-family/ChangeLog.gomp @@ -1,3 +1,10 @@ +2014-03-20 Thomas Schwinge + + * c-omp.c (check_omp_for_incr_expr, c_finish_omp_for): Update + comments. + * c-pragma.c (oacc_pragmas): Sort PRAGMA_OACC_LOOP alphabetically. + * c-pragma.h (enum pragma_kind): Likewise. + 2014-03-18 Ilmir Usmanov * c-pragma.h (enum pragma_kind): Add PRAGMA_OACC_LOOP. diff --git gcc/c-family/c-omp.c gcc/c-family/c-omp.c index 06f5712..5a1fb6d 100644 --- gcc/c-family/c-omp.c +++ gcc/c-family/c-omp.c @@ -293,7 +293,7 @@ c_finish_omp_flush (location_t loc) } -/* Check and canonicalize #pragma omp for increment expression. +/* Check and canonicalize OMP_FOR increment expression. Helper function for c_finish_omp_for. */ static tree @@ -381,7 +381,7 @@ c_omp_for_incr_canonicalize_ptr (location_t loc, tree decl, tree incr) return incr; } -/* Validate and emit code for the OpenMP directive #pragma omp for. +/* Validate and generate OMP_FOR. DECLV is a vector of iteration variables, for each collapsed loop. INITV, CONDV and INCRV are vectors containing initialization expressions, controlling predicates and increment expressions. diff --git gcc/c-family/c-pragma.c gcc/c-family/c-pragma.c index f99b087..aef4f10 100644 --- gcc/c-family/c-pragma.c +++ gcc/c-family/c-pragma.c @@ -1171,8 +1171,8 @@ struct omp_pragma_def { const char *name; unsigned int id; }; static const struct omp_pragma_def oacc_pragmas[] = { { "data", PRAGMA_OACC_DATA }, { "kernels", PRAGMA_OACC_KERNELS }, - { "parallel", PRAGMA_OACC_PARALLEL }, { "loop", PRAGMA_OACC_LOOP }, + { "parallel", PRAGMA_OACC_PARALLEL }, }; static const struct omp_pragma_def omp_pragmas[] = { { "atomic", PRAGMA_OMP_ATOMIC }, diff --git gcc/c-family/c-pragma.h gcc/c-family/c-pragma.h index f4b5b80..bb9c367 100644 --- gcc/c-family/c-pragma.h +++ gcc/c-family/c-pragma.h @@ -29,8 +29,8 @@ typedef enum pragma_kind { PRAGMA_OACC_DATA, PRAGMA_OACC_KERNELS, - PRAGMA_OACC_PARALLEL, PRAGMA_OACC_LOOP, + PRAGMA_OACC_PARALLEL, PRAGMA_OMP_ATOMIC, PRAGMA_OMP_BARRIER, PRAGMA_OMP_CANCEL, diff --git gcc/c/ChangeLog.gomp gcc/c/ChangeLog.gomp index 0358b3a..91978db 100644 --- gcc/c/ChangeLog.gomp +++ gcc/c/ChangeLog.gomp @@ -1,3 +1,11 @@ +2014-03-20 Thomas Schwinge + + * c-parser.c: Update comments. + (c_parser_oacc_loop): Move. + (c_parser_omp_for_loop): Catch some unsupported cases. + (c_parser_omp_construct) : Sort + alphabetically. + 2014-03-18 Ilmir Usmanov * c-parser.c (c_parser_oacc_loop): New function. diff --git gcc/c/c-parser.c gcc/c/c-parser.c index 734d44e..90d0035 100644 --- gcc/c/c-parser.c +++ gcc/c/c-parser.c @@ -1204,10 +1204,13 @@ static struct c_expr c_parser_expression_conv (c_parser *); static vec *c_parser_expr_list (c_parser *, bool, bool, vec **, location_t *, tree *, vec *); +static tree c_parser_oacc_loop (location_t, c_parser *, char *); static void c_parser_omp_construct (c_parser *); static void c_parser_omp_threadprivate (c_parser *); static void c_parser_omp_barrier (c_parser *); static void c_parser_omp_flush (c_parser *); +static tree c_parser_omp_for_loop (location_t, c_parser *, enum tree_code, + tree, tree *); static void c_parser_omp_taskwait (c_parser *); static void c_parser_omp_taskyield (c_parser *); static void c_parser_omp_cancel (c_parser *); @@ -4778,6 +4781,7 @@ c_parser_label (c_parser *parser) parallel-construct kernels-construct data-construct + loop-construct parallel-construct: parallel-directive structured-block @@ -4788,6 +4792,9 @@ c_parser_label (c_parser *parser) data-construct: data-directive structured-block + loop-construct: + loop-directive structured-block + OpenMP: statement: @@ -11557,8 +11564,6 @@ c_parser_oacc_data (location_t loc, c_parser *parser) return stmt; } -static tree c_parser_oacc_loop (location_t, c_parser *, char *); - /* OpenACC 2.0: # pragma acc kernels oacc-kernels-clause[optseq] new-line structured-block @@ -11611,6 +11616,33 @@ c_parser_oacc_kernels (location_t loc, c_parser *parser, char *p_name) } /* OpenACC 2.0: + # pragma acc loop oacc-loop-clause[optseq] new-line + structured-block + + LOC is the location of the #pragma token. +*/ + +#define OACC_LOOP_CLAUSE_MASK \ + (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NONE) + +static tree +c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name) +{ + tree stmt, clauses, block; + + strcat (p_name, " loop"); + + clauses = c_parser_oacc_all_clauses (parser, OACC_LOOP_CLAUSE_MASK, p_name); + + block = c_begin_compound_stmt (true); + stmt = c_parser_omp_for_loop (loc, parser, OACC_LOOP, clauses, NULL); + block = c_end_compound_stmt (loc, block, true); + add_stmt (block); + + return stmt; +} + +/* OpenACC 2.0: # pragma acc parallel oacc-parallel-clause[optseq] new-line structured-block @@ -12120,10 +12152,11 @@ c_parser_omp_flush (c_parser *parser) c_finish_omp_flush (loc); } -/* Parse the restricted form of the for statement allowed by OpenMP. +/* Parse the restricted form of loop statements allowed by OpenACC and OpenMP. The real trick here is to determine the loop control variable early so that we can push a new decl if necessary to make it private. - LOC is the location of the OMP in "#pragma omp". */ + LOC is the location of the "acc" or "omp" in "#pragma acc" or "#pragma omp", + respectively. */ static tree c_parser_omp_for_loop (location_t loc, c_parser *parser, enum tree_code code, @@ -12138,7 +12171,10 @@ c_parser_omp_for_loop (location_t loc, c_parser *parser, enum tree_code code, for (cl = clauses; cl; cl = OMP_CLAUSE_CHAIN (cl)) if (OMP_CLAUSE_CODE (cl) == OMP_CLAUSE_COLLAPSE) + { + gcc_assert (code != OACC_LOOP); collapse = tree_to_shwi (OMP_CLAUSE_COLLAPSE_EXPR (cl)); + } gcc_assert (collapse >= 1); @@ -12369,6 +12405,7 @@ c_parser_omp_for_loop (location_t loc, c_parser *parser, enum tree_code code, if (cclauses != NULL && cclauses[C_OMP_CLAUSE_SPLIT_PARALLEL] != NULL) { + gcc_assert (code != OACC_LOOP); tree *c; for (c = &cclauses[C_OMP_CLAUSE_SPLIT_PARALLEL]; *c ; ) if (OMP_CLAUSE_CODE (*c) != OMP_CLAUSE_FIRSTPRIVATE @@ -12433,33 +12470,6 @@ omp_split_clauses (location_t loc, enum tree_code code, cclauses[i] = c_finish_omp_clauses (cclauses[i]); } -/* OpenACC 2.0: - # pragma acc loop oacc-loop-clause[optseq] new-line - structured-block - - LOC is the location of the #pragma token. -*/ - -#define OACC_LOOP_CLAUSE_MASK \ - (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NONE) - -static tree -c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name) -{ - tree block, clauses, ret; - - strcat (p_name, " loop"); - - clauses = c_parser_oacc_all_clauses (parser, OACC_LOOP_CLAUSE_MASK, p_name); - - block = c_begin_compound_stmt (true); - ret = c_parser_omp_for_loop (loc, parser, OACC_LOOP, clauses, NULL); - block = c_end_compound_stmt (loc, block, true); - add_stmt (block); - - return ret; -} - /* OpenMP 4.0: #pragma omp simd simd-clause[optseq] new-line for-loop @@ -13977,14 +13987,14 @@ c_parser_omp_construct (c_parser *parser) strcpy (p_name, "#pragma acc"); stmt = c_parser_oacc_kernels (loc, parser, p_name); break; - case PRAGMA_OACC_PARALLEL: - strcpy (p_name, "#pragma acc"); - stmt = c_parser_oacc_parallel (loc, parser, p_name); - break; case PRAGMA_OACC_LOOP: strcpy (p_name, "#pragma acc"); stmt = c_parser_oacc_loop (loc, parser, p_name); break; + case PRAGMA_OACC_PARALLEL: + strcpy (p_name, "#pragma acc"); + stmt = c_parser_oacc_parallel (loc, parser, p_name); + break; case PRAGMA_OMP_ATOMIC: c_parser_omp_atomic (loc, parser); return; diff --git gcc/doc/generic.texi gcc/doc/generic.texi index 0a77a86..7780fe8 100644 --- gcc/doc/generic.texi +++ gcc/doc/generic.texi @@ -2054,8 +2054,8 @@ edge. Rethrowing the exception is represented using @code{RESX_EXPR}. @tindex OACC_PARALLEL @tindex OACC_KERNELS @tindex OACC_DATA -@tindex OACC_LOOP @tindex OACC_HOST_DATA +@tindex OACC_LOOP @tindex OACC_DECLARE @tindex OACC_UPDATE @tindex OACC_ENTER_DATA @@ -2091,14 +2091,16 @@ Represents @code{#pragma acc kernels [clause1 @dots{} clauseN]}. Represents @code{#pragma acc data [clause1 @dots{} clauseN]}. -@item OACC_LOOP - -Represents @code{#pragma acc loop [clause1 @dots{} clauseN]}. - @item OACC_HOST_DATA Represents @code{#pragma acc host_data [clause1 @dots{} clauseN]}. +@item OACC_LOOP + +Represents @code{#pragma acc loop [clause1 @dots{} clauseN]}. + +See the description of the @code{OMP_FOR} code. + @item OACC_DECLARE Represents @code{#pragma acc declare [clause1 @dots{} clauseN]}. @@ -2150,8 +2152,8 @@ variables. @item OMP_FOR -Represents @code{#pragma omp for [clause1 @dots{} clauseN]}. It -has 5 operands: +Represents @code{#pragma omp for [clause1 @dots{} clauseN]}. It has +six operands: Operand @code{OMP_FOR_BODY} contains the loop body. @@ -2241,10 +2243,9 @@ building code (@code{omp-low.c}). @item OMP_CONTINUE Similarly, this instruction does not represent an OpenMP -directive, it is used by @code{OMP_FOR} and +directive, it is used by @code{OACC_LOOP}, @code{OMP_FOR} as well as @code{OMP_SECTIONS} to mark the place where the code needs to -loop to the next iteration (in the case of @code{OMP_FOR}) or -the next section (in the case of @code{OMP_SECTIONS}). +loop to the next iteration, or the next section, respectively. In some cases, @code{OMP_CONTINUE} is placed right before @code{OMP_RETURN}. But if there are cleanups that need to diff --git gcc/doc/gimple.texi gcc/doc/gimple.texi index 91748a6..fd6feae 100644 --- gcc/doc/gimple.texi +++ gcc/doc/gimple.texi @@ -1679,9 +1679,8 @@ Set @code{NAME} to be the name associated with @code{OMP} critical statement @co tree clauses, tree index, tree initial, tree final, tree incr, @ gimple_seq pre_body, enum tree_code omp_for_cond) Build a @code{GIMPLE_OMP_FOR} statement. @code{BODY} is sequence of statements -inside the for loop. @code{CLAUSES}, are any of the @code{OMP} loop -construct's clauses: private, firstprivate, lastprivate, -reductions, ordered, schedule, and nowait. @code{PRE_BODY} is the +inside the for loop. @code{CLAUSES}, are any of the loop +construct's clauses. @code{PRE_BODY} is the sequence of statements that are loop invariant. @code{INDEX} is the index variable. @code{INITIAL} is the initial value of @code{INDEX}. @code{FINAL} is final value of @code{INDEX}. OMP_FOR_COND is the predicate used to diff --git gcc/gimple-pretty-print.c gcc/gimple-pretty-print.c index c62c517..f251060 100644 --- gcc/gimple-pretty-print.c +++ gcc/gimple-pretty-print.c @@ -1116,15 +1116,18 @@ dump_gimple_omp_for (pretty_printer *buffer, gimple gs, int spc, int flags) case GF_OMP_FOR_KIND_FOR: kind = ""; break; + case GF_OMP_FOR_KIND_DISTRIBUTE: + kind = " distribute"; + break; + case GF_OMP_FOR_KIND_OACC_LOOP: + kind = " oacc_loop"; + break; case GF_OMP_FOR_KIND_SIMD: kind = " simd"; break; case GF_OMP_FOR_KIND_CILKSIMD: kind = " cilksimd"; break; - case GF_OMP_FOR_KIND_DISTRIBUTE: - kind = " distribute"; - break; default: gcc_unreachable (); } @@ -1150,15 +1153,18 @@ dump_gimple_omp_for (pretty_printer *buffer, gimple gs, int spc, int flags) case GF_OMP_FOR_KIND_FOR: pp_string (buffer, "#pragma omp for"); break; + case GF_OMP_FOR_KIND_DISTRIBUTE: + pp_string (buffer, "#pragma omp distribute"); + break; + case GF_OMP_FOR_KIND_OACC_LOOP: + pp_string (buffer, "#pragma acc loop"); + break; case GF_OMP_FOR_KIND_SIMD: pp_string (buffer, "#pragma omp simd"); break; case GF_OMP_FOR_KIND_CILKSIMD: pp_string (buffer, "#pragma simd"); break; - case GF_OMP_FOR_KIND_DISTRIBUTE: - pp_string (buffer, "#pragma omp distribute"); - break; default: gcc_unreachable (); } diff --git gcc/gimple.c gcc/gimple.c index 1862de2..6580d10 100644 --- gcc/gimple.c +++ gcc/gimple.c @@ -853,8 +853,7 @@ gimple_build_omp_critical (gimple_seq body, tree name) BODY is sequence of statements inside the for loop. KIND is the `for' variant. - CLAUSES, are any of the OMP loop construct's clauses: private, firstprivate, - lastprivate, reductions, ordered, schedule, and nowait. + CLAUSES, are any of the loop construct's clauses. COLLAPSE is the collapse count. PRE_BODY is the sequence of statements that are loop invariant. */ @@ -1694,6 +1693,7 @@ gimple_copy (gimple stmt) gcc_unreachable (); case GIMPLE_OMP_FOR: + gcc_assert (!is_gimple_omp_oacc_specifically (stmt)); new_seq = gimple_seq_copy (gimple_omp_for_pre_body (stmt)); gimple_omp_for_set_pre_body (copy, new_seq); t = unshare_expr (gimple_omp_for_clauses (stmt)); diff --git gcc/gimple.def gcc/gimple.def index c9756b7..e2e912c 100644 --- gcc/gimple.def +++ gcc/gimple.def @@ -267,6 +267,9 @@ DEFGSCODE(GIMPLE_OMP_CRITICAL, "gimple_omp_critical", GSS_OMP_CRITICAL) for (INDEX = INITIAL; INDEX COND FINAL; INDEX {+=,-=} INCR) BODY + Likewise for: + #pragma acc loop [clause1 ... clauseN] + BODY is the loop body. CLAUSES is the list of clauses. @@ -293,7 +296,7 @@ DEFGSCODE(GIMPLE_OMP_CRITICAL, "gimple_omp_critical", GSS_OMP_CRITICAL) INITIAL, FINAL and INCR are required to be loop invariant integer expressions that are evaluated without any synchronization. The evaluation order, frequency of evaluation and side-effects are - unspecified by the standard. */ + unspecified by the standards. */ DEFGSCODE(GIMPLE_OMP_FOR, "gimple_omp_for", GSS_OMP_FOR) /* GIMPLE_OMP_MASTER represents #pragma omp master. diff --git gcc/gimple.h gcc/gimple.h index 34a0bdb..f059789 100644 --- gcc/gimple.h +++ gcc/gimple.h @@ -91,15 +91,16 @@ enum gf_mask { GF_CALL_ALLOCA_FOR_VAR = 1 << 5, GF_CALL_INTERNAL = 1 << 6, GF_OMP_PARALLEL_COMBINED = 1 << 0, - GF_OMP_FOR_KIND_MASK = (1 << 2) - 1, + GF_OMP_FOR_KIND_MASK = (1 << 3) - 1, GF_OMP_FOR_KIND_FOR = 0, GF_OMP_FOR_KIND_DISTRIBUTE = 1, + GF_OMP_FOR_KIND_OACC_LOOP = 2, /* Flag for SIMD variants of OMP_FOR kinds. */ - GF_OMP_FOR_SIMD = 1 << 1, + GF_OMP_FOR_SIMD = 1 << 2, GF_OMP_FOR_KIND_SIMD = GF_OMP_FOR_SIMD | 0, GF_OMP_FOR_KIND_CILKSIMD = GF_OMP_FOR_SIMD | 1, - GF_OMP_FOR_COMBINED = 1 << 2, - GF_OMP_FOR_COMBINED_INTO = 1 << 3, + GF_OMP_FOR_COMBINED = 1 << 3, + GF_OMP_FOR_COMBINED_INTO = 1 << 4, GF_OMP_TARGET_KIND_MASK = (1 << 2) - 1, GF_OMP_TARGET_KIND_REGION = 0, GF_OMP_TARGET_KIND_DATA = 1, @@ -4518,7 +4519,7 @@ gimple_omp_critical_set_name (gimple gs, tree name) } -/* Return the kind of OMP for statemement. */ +/* Return the kind of the OMP_FOR statemement G. */ static inline int gimple_omp_for_kind (const_gimple g) @@ -4528,7 +4529,7 @@ gimple_omp_for_kind (const_gimple g) } -/* Set the OMP for kind. */ +/* Set the kind of the OMP_FOR statement G. */ static inline void gimple_omp_for_set_kind (gimple g, int kind) @@ -4539,7 +4540,7 @@ gimple_omp_for_set_kind (gimple g, int kind) } -/* Return true if OMP for statement G has the +/* Return true if OMP_FOR statement G has the GF_OMP_FOR_COMBINED flag set. */ static inline bool @@ -4550,8 +4551,8 @@ gimple_omp_for_combined_p (const_gimple g) } -/* Set the GF_OMP_FOR_COMBINED field in G depending on the boolean - value of COMBINED_P. */ +/* Set the GF_OMP_FOR_COMBINED field in the OMP_FOR statement G depending on + the boolean value of COMBINED_P. */ static inline void gimple_omp_for_set_combined_p (gimple g, bool combined_p) @@ -4564,7 +4565,7 @@ gimple_omp_for_set_combined_p (gimple g, bool combined_p) } -/* Return true if OMP for statement G has the +/* Return true if the OMP_FOR statement G has the GF_OMP_FOR_COMBINED_INTO flag set. */ static inline bool @@ -4575,8 +4576,8 @@ gimple_omp_for_combined_into_p (const_gimple g) } -/* Set the GF_OMP_FOR_COMBINED_INTO field in G depending on the boolean - value of COMBINED_P. */ +/* Set the GF_OMP_FOR_COMBINED_INTO field in the OMP_FOR statement G depending + on the boolean value of COMBINED_P. */ static inline void gimple_omp_for_set_combined_into_p (gimple g, bool combined_p) @@ -4589,7 +4590,7 @@ gimple_omp_for_set_combined_into_p (gimple g, bool combined_p) } -/* Return the clauses associated with OMP_FOR GS. */ +/* Return the clauses associated with the OMP_FOR statement GS. */ static inline tree gimple_omp_for_clauses (const_gimple gs) @@ -4600,7 +4601,8 @@ gimple_omp_for_clauses (const_gimple gs) } -/* Return a pointer to the OMP_FOR GS. */ +/* Return a pointer to the clauses associated with the OMP_FOR statement + GS. */ static inline tree * gimple_omp_for_clauses_ptr (gimple gs) @@ -4611,7 +4613,8 @@ gimple_omp_for_clauses_ptr (gimple gs) } -/* Set CLAUSES to be the list of clauses associated with OMP_FOR GS. */ +/* Set CLAUSES to be the list of clauses associated with the OMP_FOR statement + GS. */ static inline void gimple_omp_for_set_clauses (gimple gs, tree clauses) @@ -4622,7 +4625,7 @@ gimple_omp_for_set_clauses (gimple gs, tree clauses) } -/* Get the collapse count of OMP_FOR GS. */ +/* Get the collapse count of the OMP_FOR statement GS. */ static inline size_t gimple_omp_for_collapse (gimple gs) @@ -4633,7 +4636,7 @@ gimple_omp_for_collapse (gimple gs) } -/* Return the index variable for OMP_FOR GS. */ +/* Return the index variable for the OMP_FOR statement GS. */ static inline tree gimple_omp_for_index (const_gimple gs, size_t i) @@ -4645,7 +4648,7 @@ gimple_omp_for_index (const_gimple gs, size_t i) } -/* Return a pointer to the index variable for OMP_FOR GS. */ +/* Return a pointer to the index variable for the OMP_FOR statement GS. */ static inline tree * gimple_omp_for_index_ptr (gimple gs, size_t i) @@ -4657,7 +4660,7 @@ gimple_omp_for_index_ptr (gimple gs, size_t i) } -/* Set INDEX to be the index variable for OMP_FOR GS. */ +/* Set INDEX to be the index variable for the OMP_FOR statement GS. */ static inline void gimple_omp_for_set_index (gimple gs, size_t i, tree index) @@ -4669,7 +4672,7 @@ gimple_omp_for_set_index (gimple gs, size_t i, tree index) } -/* Return the initial value for OMP_FOR GS. */ +/* Return the initial value for the OMP_FOR statement GS. */ static inline tree gimple_omp_for_initial (const_gimple gs, size_t i) @@ -4681,7 +4684,7 @@ gimple_omp_for_initial (const_gimple gs, size_t i) } -/* Return a pointer to the initial value for OMP_FOR GS. */ +/* Return a pointer to the initial value for the OMP_FOR statement GS. */ static inline tree * gimple_omp_for_initial_ptr (gimple gs, size_t i) @@ -4693,7 +4696,7 @@ gimple_omp_for_initial_ptr (gimple gs, size_t i) } -/* Set INITIAL to be the initial value for OMP_FOR GS. */ +/* Set INITIAL to be the initial value for the OMP_FOR statement GS. */ static inline void gimple_omp_for_set_initial (gimple gs, size_t i, tree initial) @@ -4705,7 +4708,7 @@ gimple_omp_for_set_initial (gimple gs, size_t i, tree initial) } -/* Return the final value for OMP_FOR GS. */ +/* Return the final value for the OMP_FOR statement GS. */ static inline tree gimple_omp_for_final (const_gimple gs, size_t i) @@ -4717,7 +4720,7 @@ gimple_omp_for_final (const_gimple gs, size_t i) } -/* Return a pointer to the final value for OMP_FOR GS. */ +/* Return a pointer to the final value for the OMP_FOR statement GS. */ static inline tree * gimple_omp_for_final_ptr (gimple gs, size_t i) @@ -4729,7 +4732,7 @@ gimple_omp_for_final_ptr (gimple gs, size_t i) } -/* Set FINAL to be the final value for OMP_FOR GS. */ +/* Set FINAL to be the final value for the OMP_FOR statement GS. */ static inline void gimple_omp_for_set_final (gimple gs, size_t i, tree final) @@ -4741,7 +4744,32 @@ gimple_omp_for_set_final (gimple gs, size_t i, tree final) } -/* Return the increment value for OMP_FOR GS. */ +/* Set COND to be the condition code for the OMP_FOR statement GS. */ + +static inline void +gimple_omp_for_set_cond (gimple gs, size_t i, enum tree_code cond) +{ + gimple_statement_omp_for *omp_for_stmt = + as_a (gs); + gcc_gimple_checking_assert (TREE_CODE_CLASS (cond) == tcc_comparison + && i < omp_for_stmt->collapse); + omp_for_stmt->iter[i].cond = cond; +} + + +/* Return the condition code associated with the OMP_FOR statement GS. */ + +static inline enum tree_code +gimple_omp_for_cond (const_gimple gs, size_t i) +{ + const gimple_statement_omp_for *omp_for_stmt = + as_a (gs); + gcc_gimple_checking_assert (i < omp_for_stmt->collapse); + return omp_for_stmt->iter[i].cond; +} + + +/* Return the increment value for the OMP_FOR statement GS. */ static inline tree gimple_omp_for_incr (const_gimple gs, size_t i) @@ -4753,7 +4781,7 @@ gimple_omp_for_incr (const_gimple gs, size_t i) } -/* Return a pointer to the increment value for OMP_FOR GS. */ +/* Return a pointer to the increment value for the OMP_FOR statement GS. */ static inline tree * gimple_omp_for_incr_ptr (gimple gs, size_t i) @@ -4765,7 +4793,7 @@ gimple_omp_for_incr_ptr (gimple gs, size_t i) } -/* Set INCR to be the increment value for OMP_FOR GS. */ +/* Set INCR to be the increment value for the OMP_FOR statement GS. */ static inline void gimple_omp_for_set_incr (gimple gs, size_t i, tree incr) @@ -5470,31 +5498,6 @@ gimple_omp_sections_set_control (gimple gs, tree control) } -/* Set COND to be the condition code for OMP_FOR GS. */ - -static inline void -gimple_omp_for_set_cond (gimple gs, size_t i, enum tree_code cond) -{ - gimple_statement_omp_for *omp_for_stmt = - as_a (gs); - gcc_gimple_checking_assert (TREE_CODE_CLASS (cond) == tcc_comparison - && i < omp_for_stmt->collapse); - omp_for_stmt->iter[i].cond = cond; -} - - -/* Return the condition code associated with OMP_FOR GS. */ - -static inline enum tree_code -gimple_omp_for_cond (const_gimple gs, size_t i) -{ - const gimple_statement_omp_for *omp_for_stmt = - as_a (gs); - gcc_gimple_checking_assert (i < omp_for_stmt->collapse); - return omp_for_stmt->iter[i].cond; -} - - /* Set the value being stored in an atomic store. */ static inline void @@ -5811,6 +5814,14 @@ is_gimple_omp_oacc_specifically (const_gimple stmt) case GIMPLE_OACC_KERNELS: case GIMPLE_OACC_PARALLEL: return true; + case GIMPLE_OMP_FOR: + switch (gimple_omp_for_kind (stmt)) + { + case GF_OMP_FOR_KIND_OACC_LOOP: + return true; + default: + return false; + } case GIMPLE_OMP_TARGET: switch (gimple_omp_target_kind (stmt)) { diff --git gcc/gimplify.c gcc/gimplify.c index 0985bb2..9788f4c 100644 --- gcc/gimplify.c +++ gcc/gimplify.c @@ -4363,8 +4363,8 @@ is_gimple_stmt (tree t) case OMP_FOR: case OMP_SIMD: case CILK_SIMD: - case OACC_LOOP: case OMP_DISTRIBUTE: + case OACC_LOOP: case OMP_SECTIONS: case OMP_SECTION: case OMP_SINGLE: @@ -6683,14 +6683,36 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p) gimple_seq for_body, for_pre_body; int i; bool simd; + enum gimplify_omp_var_data govd_private; + enum omp_region_type ort; bitmap has_decl_expr = NULL; orig_for_stmt = for_stmt = *expr_p; - simd = TREE_CODE (for_stmt) == OMP_SIMD - || TREE_CODE (for_stmt) == CILK_SIMD; - gimplify_scan_omp_clauses (&OMP_FOR_CLAUSES (for_stmt), pre_p, - simd ? ORT_SIMD : ORT_WORKSHARE); + switch (TREE_CODE (for_stmt)) + { + case OMP_FOR: + case OMP_DISTRIBUTE: + simd = false; + govd_private = GOVD_PRIVATE; + ort = ORT_WORKSHARE; + break; + case OACC_LOOP: + simd = false; + govd_private = /* TODO */ GOVD_LOCAL; + ort = /* TODO */ ORT_WORKSHARE; + break; + case OMP_SIMD: + case CILK_SIMD: + simd = true; + govd_private = GOVD_PRIVATE; + ort = ORT_SIMD; + break; + default: + gcc_unreachable (); + } + + gimplify_scan_omp_clauses (&OMP_FOR_CLAUSES (for_stmt), pre_p, ort); /* Handle OMP_FOR_INIT. */ for_pre_body = NULL; @@ -6722,6 +6744,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p) if (OMP_FOR_INIT (for_stmt) == NULL_TREE) { + gcc_assert (TREE_CODE (for_stmt) != OACC_LOOP); for_stmt = walk_tree (&OMP_FOR_BODY (for_stmt), find_combined_omp_for, NULL, NULL); gcc_assert (for_stmt != NULL_TREE); @@ -6742,7 +6765,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p) gcc_assert (INTEGRAL_TYPE_P (TREE_TYPE (decl)) || POINTER_TYPE_P (TREE_TYPE (decl))); - /* Make sure the iteration variable is private. */ + /* Make sure the iteration variable is some kind of private. */ tree c = NULL_TREE; if (orig_for_stmt != for_stmt) /* Do this only on innermost construct for combined ones. */; @@ -6768,6 +6791,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p) } else { + gcc_assert (govd_private == GOVD_PRIVATE); bool lastprivate = (!has_decl_expr || !bitmap_bit_p (has_decl_expr, DECL_UID (decl))); @@ -6785,7 +6809,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p) else if (omp_is_private (gimplify_omp_ctxp, decl, simd)) omp_notice_variable (gimplify_omp_ctxp, decl, true); else - 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 @@ -6799,7 +6823,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p) gimplify_seq_add_stmt (&for_body, gimple_build_assign (decl, var)); - omp_add_variable (gimplify_omp_ctxp, var, GOVD_PRIVATE | GOVD_SEEN); + omp_add_variable (gimplify_omp_ctxp, var, govd_private | GOVD_SEEN); } else var = decl; @@ -6936,7 +6960,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p) t = TREE_VEC_ELT (OMP_FOR_INIT (for_stmt), i); decl = TREE_OPERAND (t, 0); var = create_tmp_var (TREE_TYPE (decl), get_name (decl)); - omp_add_variable (gimplify_omp_ctxp, var, GOVD_PRIVATE | GOVD_SEEN); + omp_add_variable (gimplify_omp_ctxp, var, govd_private | GOVD_SEEN); TREE_OPERAND (t, 0) = var; t = TREE_VEC_ELT (OMP_FOR_INCR (for_stmt), i); TREE_OPERAND (t, 1) = copy_node (TREE_OPERAND (t, 1)); @@ -6952,6 +6976,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p) case OMP_SIMD: kind = GF_OMP_FOR_KIND_SIMD; break; case CILK_SIMD: kind = GF_OMP_FOR_KIND_CILKSIMD; break; case OMP_DISTRIBUTE: kind = GF_OMP_FOR_KIND_DISTRIBUTE; break; + case OACC_LOOP: kind = GF_OMP_FOR_KIND_OACC_LOOP; break; default: gcc_unreachable (); } @@ -8048,7 +8073,6 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p, case OACC_EXIT_DATA: case OACC_WAIT: case OACC_CACHE: - case OACC_LOOP: sorry ("directive not yet implemented"); ret = GS_ALL_DONE; break; @@ -8067,6 +8091,7 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p, case OMP_SIMD: case CILK_SIMD: case OMP_DISTRIBUTE: + case OACC_LOOP: ret = gimplify_omp_for (expr_p, pre_p); break; diff --git gcc/omp-low.c gcc/omp-low.c index c3b3e95..13373ca 100644 --- gcc/omp-low.c +++ gcc/omp-low.c @@ -177,6 +177,8 @@ typedef struct omp_context bool cancellable; } omp_context; +/* A structure holding the elements of: + for (V = N1; V cond N2; V += STEP) [...] */ struct omp_for_data_loop { @@ -310,9 +312,9 @@ extract_omp_for_data (gimple for_stmt, struct omp_for_data *fd, else fd->loops = &fd->loop; - fd->have_nowait = distribute || simd; + fd->have_nowait = (gimple_omp_for_kind (for_stmt) != GF_OMP_FOR_KIND_FOR); fd->have_ordered = false; - fd->sched_kind = OMP_CLAUSE_SCHEDULE_STATIC; + fd->sched_kind = /* TODO: OACC_LOOP */ OMP_CLAUSE_SCHEDULE_STATIC; fd->chunk_size = NULL_TREE; collapse_iter = NULL; collapse_count = NULL; @@ -1626,7 +1628,10 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_NUM_WORKERS: case OMP_CLAUSE_VECTOR_LENGTH: if (ctx->outer) + { + gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt)); scan_omp_op (&OMP_CLAUSE_OPERAND (c, 0), ctx->outer); + } break; case OMP_CLAUSE_TO: @@ -2288,7 +2293,7 @@ scan_omp_task (gimple_stmt_iterator *gsi, omp_context *outer_ctx) } -/* Scan an OpenMP loop directive. */ +/* Scan a GIMPLE_OMP_FOR. */ static void scan_omp_for (gimple stmt, omp_context *outer_ctx) @@ -2421,6 +2426,10 @@ check_omp_nesting_restrictions (gimple stmt, omp_context *ctx) if (is_gimple_omp (stmt) && is_gimple_omp_oacc_specifically (stmt)) { + /* Regular handling of OpenACC loop constructs. */ + if (gimple_code (stmt) == GIMPLE_OMP_FOR + && gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_OACC_LOOP) + goto cont; /* No nesting of OpenACC STMT inside any OpenACC or OpenMP CTX different from an OpenACC data construct. */ for (omp_context *ctx_ = ctx; ctx_ != NULL; ctx_ = ctx_->outer) @@ -2447,6 +2456,7 @@ check_omp_nesting_restrictions (gimple stmt, omp_context *ctx) return false; } } + cont: if (ctx != NULL) { @@ -2626,6 +2636,8 @@ check_omp_nesting_restrictions (gimple stmt, omp_context *ctx) "of work-sharing, critical, ordered, master or explicit " "task region"); return false; + case GIMPLE_OACC_KERNELS: + case GIMPLE_OACC_PARALLEL: case GIMPLE_OMP_PARALLEL: return true; default: @@ -3217,8 +3229,6 @@ static void lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist, omp_context *ctx, struct omp_for_data *fd) { - gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt)); - tree c, dtor, copyin_seq, x, ptr; bool copyin_by_ref = false; bool lastprivate_firstprivate = false; @@ -3920,8 +3930,6 @@ static void lower_lastprivate_clauses (tree clauses, tree predicate, gimple_seq *stmt_list, omp_context *ctx) { - gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt)); - tree x, c, label = NULL, orig_clauses = clauses; bool par_clauses = false; tree simduid = NULL, lastlane = NULL; @@ -4057,8 +4065,6 @@ lower_lastprivate_clauses (tree clauses, tree predicate, gimple_seq *stmt_list, static void lower_reduction_clauses (tree clauses, gimple_seq *stmt_seqp, omp_context *ctx) { - gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt)); - gimple_seq sub_seq = NULL; gimple stmt; tree x, c; @@ -5849,6 +5855,8 @@ expand_omp_for_generic (struct omp_region *region, enum built_in_function next_fn, gimple inner_stmt) { + gcc_assert (gimple_omp_for_kind (fd->for_stmt) != GF_OMP_FOR_KIND_OACC_LOOP); + tree type, istart0, iend0, iend; tree t, vmain, vback, bias = NULL_TREE; basic_block entry_bb, cont_bb, exit_bb, l0_bb, l1_bb, collapse_bb; @@ -5918,6 +5926,9 @@ expand_omp_for_generic (struct omp_region *region, gcc_assert (gimple_code (gsi_stmt (gsi)) == GIMPLE_OMP_FOR); if (fd->collapse > 1) { + gcc_assert (gimple_omp_for_kind (gsi_stmt (gsi)) + != GF_OMP_FOR_KIND_OACC_LOOP); + int first_zero_iter = -1; basic_block zero_iter_bb = NULL, l2_dom_bb = NULL; @@ -5946,6 +5957,9 @@ expand_omp_for_generic (struct omp_region *region, } if (in_combined_parallel) { + gcc_assert (gimple_omp_for_kind (gsi_stmt (gsi)) + != GF_OMP_FOR_KIND_OACC_LOOP); + /* In a combined parallel loop, emit a call to GOMP_loop_foo_next. */ t = build_call_expr (builtin_decl_explicit (next_fn), 2, @@ -5964,6 +5978,9 @@ expand_omp_for_generic (struct omp_region *region, t0 = fd->loop.n1; if (gimple_omp_for_combined_into_p (fd->for_stmt)) { + gcc_assert (gimple_omp_for_kind (gsi_stmt (gsi)) + != GF_OMP_FOR_KIND_OACC_LOOP); + tree innerc = find_omp_clause (gimple_omp_for_clauses (fd->for_stmt), OMP_CLAUSE__LOOPTEMP_); gcc_assert (innerc); @@ -6276,12 +6293,14 @@ expand_omp_for_static_nochunk (struct omp_region *region, gimple_stmt_iterator gsi; gimple stmt; edge ep; - enum built_in_function get_num_threads = BUILT_IN_OMP_GET_NUM_THREADS; - enum built_in_function get_thread_num = BUILT_IN_OMP_GET_THREAD_NUM; bool broken_loop = region->cont == NULL; tree *counts = NULL; tree n1, n2, step; + gcc_assert ((gimple_omp_for_kind (fd->for_stmt) + != GF_OMP_FOR_KIND_OACC_LOOP) + || !inner_stmt); + itype = type = TREE_TYPE (fd->loop.v); if (POINTER_TYPE_P (type)) itype = signed_type_for (type); @@ -6305,14 +6324,11 @@ expand_omp_for_static_nochunk (struct omp_region *region, gsi = gsi_last_bb (entry_bb); gcc_assert (gimple_code (gsi_stmt (gsi)) == GIMPLE_OMP_FOR); - if (gimple_omp_for_kind (fd->for_stmt) == GF_OMP_FOR_KIND_DISTRIBUTE) - { - get_num_threads = BUILT_IN_OMP_GET_NUM_TEAMS; - get_thread_num = BUILT_IN_OMP_GET_TEAM_NUM; - } - if (fd->collapse > 1) { + gcc_assert (gimple_omp_for_kind (fd->for_stmt) + != GF_OMP_FOR_KIND_OACC_LOOP); + int first_zero_iter = -1; basic_block l2_dom_bb = NULL; @@ -6323,7 +6339,12 @@ expand_omp_for_static_nochunk (struct omp_region *region, t = NULL_TREE; } else if (gimple_omp_for_combined_into_p (fd->for_stmt)) + { + gcc_assert (gimple_omp_for_kind (fd->for_stmt) + != GF_OMP_FOR_KIND_OACC_LOOP); + t = integer_one_node; + } else t = fold_binary (fd->loop.cond_code, boolean_type_node, fold_convert (type, fd->loop.n1), @@ -6357,6 +6378,9 @@ expand_omp_for_static_nochunk (struct omp_region *region, ep->probability = REG_BR_PROB_BASE / 2000 - 1; if (gimple_in_ssa_p (cfun)) { + gcc_assert (gimple_omp_for_kind (fd->for_stmt) + != GF_OMP_FOR_KIND_OACC_LOOP); + int dest_idx = find_edge (entry_bb, fin_bb)->dest_idx; for (gsi = gsi_start_phis (fin_bb); !gsi_end_p (gsi); gsi_next (&gsi)) @@ -6369,14 +6393,32 @@ expand_omp_for_static_nochunk (struct omp_region *region, gsi = gsi_last_bb (entry_bb); } - t = build_call_expr (builtin_decl_explicit (get_num_threads), 0); - t = fold_convert (itype, t); - nthreads = force_gimple_operand_gsi (&gsi, t, true, NULL_TREE, + switch (gimple_omp_for_kind (fd->for_stmt)) + { + case GF_OMP_FOR_KIND_FOR: + nthreads = builtin_decl_explicit (BUILT_IN_OMP_GET_NUM_THREADS); + nthreads = build_call_expr (nthreads, 0); + threadid = builtin_decl_explicit (BUILT_IN_OMP_GET_THREAD_NUM); + threadid = build_call_expr (threadid, 0); + break; + case GF_OMP_FOR_KIND_DISTRIBUTE: + nthreads = builtin_decl_explicit (BUILT_IN_OMP_GET_NUM_TEAMS); + nthreads = build_call_expr (nthreads, 0); + threadid = builtin_decl_explicit (BUILT_IN_OMP_GET_TEAM_NUM); + threadid = build_call_expr (threadid, 0); + break; + case GF_OMP_FOR_KIND_OACC_LOOP: + nthreads = integer_one_node; + threadid = integer_zero_node; + break; + default: + gcc_unreachable (); + } + nthreads = fold_convert (itype, nthreads); + nthreads = force_gimple_operand_gsi (&gsi, nthreads, true, NULL_TREE, true, GSI_SAME_STMT); - - t = build_call_expr (builtin_decl_explicit (get_thread_num), 0); - t = fold_convert (itype, t); - threadid = force_gimple_operand_gsi (&gsi, t, true, NULL_TREE, + threadid = fold_convert (itype, threadid); + threadid = force_gimple_operand_gsi (&gsi, threadid, true, NULL_TREE, true, GSI_SAME_STMT); n1 = fd->loop.n1; @@ -6384,6 +6426,9 @@ expand_omp_for_static_nochunk (struct omp_region *region, step = fd->loop.step; if (gimple_omp_for_combined_into_p (fd->for_stmt)) { + gcc_assert (gimple_omp_for_kind (fd->for_stmt) + != GF_OMP_FOR_KIND_OACC_LOOP); + tree innerc = find_omp_clause (gimple_omp_for_clauses (fd->for_stmt), OMP_CLAUSE__LOOPTEMP_); gcc_assert (innerc); @@ -6462,6 +6507,9 @@ expand_omp_for_static_nochunk (struct omp_region *region, if (gimple_omp_for_combined_p (fd->for_stmt)) { + gcc_assert (gimple_omp_for_kind (fd->for_stmt) + != GF_OMP_FOR_KIND_OACC_LOOP); + tree clauses = gimple_code (inner_stmt) == GIMPLE_OMP_PARALLEL ? gimple_omp_parallel_clauses (inner_stmt) : gimple_omp_for_clauses (inner_stmt); @@ -6502,7 +6550,12 @@ expand_omp_for_static_nochunk (struct omp_region *region, gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING); } if (fd->collapse > 1) + { + gcc_assert (gimple_omp_for_kind (fd->for_stmt) + != GF_OMP_FOR_KIND_OACC_LOOP); + expand_omp_for_init_vars (fd, &gsi, counts, inner_stmt, startvar); + } if (!broken_loop) { @@ -6537,13 +6590,21 @@ expand_omp_for_static_nochunk (struct omp_region *region, gsi_remove (&gsi, true); if (fd->collapse > 1 && !gimple_omp_for_combined_p (fd->for_stmt)) + { + gcc_assert (gimple_omp_for_kind (fd->for_stmt) + != GF_OMP_FOR_KIND_OACC_LOOP); + collapse_bb = extract_omp_for_update_vars (fd, cont_bb, body_bb); + } } /* Replace the GIMPLE_OMP_RETURN with a barrier, or nothing. */ gsi = gsi_last_bb (exit_bb); if (!gimple_omp_return_nowait_p (gsi_stmt (gsi))) { + gcc_assert (gimple_omp_for_kind (fd->for_stmt) + != GF_OMP_FOR_KIND_OACC_LOOP); + t = gimple_omp_return_lhs (gsi_stmt (gsi)); gsi_insert_after (&gsi, build_omp_barrier (t), GSI_SAME_STMT); } @@ -6563,11 +6624,17 @@ expand_omp_for_static_nochunk (struct omp_region *region, ep = find_edge (cont_bb, body_bb); if (gimple_omp_for_combined_p (fd->for_stmt)) { + gcc_assert (gimple_omp_for_kind (fd->for_stmt) + != GF_OMP_FOR_KIND_OACC_LOOP); + remove_edge (ep); ep = NULL; } else if (fd->collapse > 1) { + gcc_assert (gimple_omp_for_kind (fd->for_stmt) + != GF_OMP_FOR_KIND_OACC_LOOP); + remove_edge (ep); ep = make_edge (cont_bb, collapse_bb, EDGE_TRUE_VALUE); } @@ -6639,6 +6706,8 @@ static void expand_omp_for_static_chunk (struct omp_region *region, struct omp_for_data *fd, gimple inner_stmt) { + gcc_assert (gimple_omp_for_kind (fd->for_stmt) != GF_OMP_FOR_KIND_OACC_LOOP); + tree n, s0, e0, e, t; tree trip_var, trip_init, trip_main, trip_back, nthreads, threadid; tree type, itype, v_main, v_back, v_extra; @@ -6647,8 +6716,6 @@ expand_omp_for_static_chunk (struct omp_region *region, gimple_stmt_iterator si; gimple stmt; edge se; - enum built_in_function get_num_threads = BUILT_IN_OMP_GET_NUM_THREADS; - enum built_in_function get_thread_num = BUILT_IN_OMP_GET_THREAD_NUM; bool broken_loop = region->cont == NULL; tree *counts = NULL; tree n1, n2, step; @@ -6680,12 +6747,6 @@ expand_omp_for_static_chunk (struct omp_region *region, si = gsi_last_bb (entry_bb); gcc_assert (gimple_code (gsi_stmt (si)) == GIMPLE_OMP_FOR); - if (gimple_omp_for_kind (fd->for_stmt) == GF_OMP_FOR_KIND_DISTRIBUTE) - { - get_num_threads = BUILT_IN_OMP_GET_NUM_TEAMS; - get_thread_num = BUILT_IN_OMP_GET_TEAM_NUM; - } - if (fd->collapse > 1) { int first_zero_iter = -1; @@ -6744,14 +6805,28 @@ expand_omp_for_static_chunk (struct omp_region *region, si = gsi_last_bb (entry_bb); } - t = build_call_expr (builtin_decl_explicit (get_num_threads), 0); - t = fold_convert (itype, t); - nthreads = force_gimple_operand_gsi (&si, t, true, NULL_TREE, + switch (gimple_omp_for_kind (fd->for_stmt)) + { + case GF_OMP_FOR_KIND_FOR: + nthreads = builtin_decl_explicit (BUILT_IN_OMP_GET_NUM_THREADS); + nthreads = build_call_expr (nthreads, 0); + threadid = builtin_decl_explicit (BUILT_IN_OMP_GET_THREAD_NUM); + threadid = build_call_expr (threadid, 0); + break; + case GF_OMP_FOR_KIND_DISTRIBUTE: + nthreads = builtin_decl_explicit (BUILT_IN_OMP_GET_NUM_TEAMS); + nthreads = build_call_expr (nthreads, 0); + threadid = builtin_decl_explicit (BUILT_IN_OMP_GET_TEAM_NUM); + threadid = build_call_expr (threadid, 0); + break; + default: + gcc_unreachable (); + } + nthreads = fold_convert (itype, nthreads); + nthreads = force_gimple_operand_gsi (&si, nthreads, true, NULL_TREE, true, GSI_SAME_STMT); - - t = build_call_expr (builtin_decl_explicit (get_thread_num), 0); - t = fold_convert (itype, t); - threadid = force_gimple_operand_gsi (&si, t, true, NULL_TREE, + threadid = fold_convert (itype, threadid); + threadid = force_gimple_operand_gsi (&si, threadid, true, NULL_TREE, true, GSI_SAME_STMT); n1 = fd->loop.n1; @@ -9211,8 +9286,6 @@ lower_oacc_offload (gimple_stmt_iterator *gsi_p, omp_context *ctx) static void maybe_add_implicit_barrier_cancel (omp_context *ctx, gimple_seq *body) { - gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt)); - gimple omp_return = gimple_seq_last_stmt (*body); gcc_assert (gimple_code (omp_return) == GIMPLE_OMP_RETURN); if (gimple_omp_return_nowait_p (omp_return)) @@ -9792,6 +9865,8 @@ lower_omp_for (gimple_stmt_iterator *gsi_p, omp_context *ctx) if (gimple_omp_for_combined_into_p (stmt)) { + gcc_assert (gimple_omp_for_kind (stmt) != GF_OMP_FOR_KIND_OACC_LOOP); + extract_omp_for_data (stmt, &fd, NULL); fdp = &fd; diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp index 13e99d5..78882c0 100644 --- gcc/testsuite/ChangeLog.gomp +++ gcc/testsuite/ChangeLog.gomp @@ -1,5 +1,12 @@ 2014-03-20 Thomas Schwinge + * c-c++-common/goacc-gomp/nesting-1.c: New file. + * c-c++-common/goacc-gomp/nesting-fail-1.c: Extend. + * c-c++-common/goacc/clauses-fail.c: Likewise. + * c-c++-common/goacc/nesting-1.c: Likewise. + * gcc.dg/goacc/sb-1.c: Likewise. + * gcc.dg/goacc/sb-3.c: New file. + * c-c++-common/goacc/nesting-1.c: New file. * c-c++-common/goacc/nesting-data-1.c: Likewise. * c-c++-common/goacc/nesting-fail-1.c: Update. diff --git gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c new file mode 100644 index 0000000..df45bcf --- /dev/null +++ gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c @@ -0,0 +1,12 @@ +void +f_omp_parallel (void) +{ +#pragma omp parallel + { + int i; + +#pragma acc loop + for (i = 0; i < 2; ++i) + ; + } +} diff --git gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c index 14103a6..871fab3 100644 --- gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c +++ gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c @@ -24,6 +24,9 @@ f_omp (void) ; #pragma acc data /* { dg-error "may not be nested" } */ ; +#pragma acc loop /* { dg-error "may not be closely nested" } */ + for (i = 0; i < 2; ++i) + ; } #pragma omp sections @@ -42,6 +45,12 @@ f_omp (void) #pragma acc data /* { dg-error "may not be nested" } */ ; } +#pragma omp section + { +#pragma acc loop /* { dg-error "may not be closely nested" } */ + for (i = 0; i < 2; ++i) + ; + } } #pragma omp single @@ -52,6 +61,9 @@ f_omp (void) ; #pragma acc data /* { dg-error "may not be nested" } */ ; +#pragma acc loop /* { dg-error "may not be closely nested" } */ + for (i = 0; i < 2; ++i) + ; } #pragma omp task @@ -62,6 +74,9 @@ f_omp (void) ; #pragma acc data /* { dg-error "may not be nested" } */ ; +#pragma acc loop /* { dg-error "may not be closely nested" } */ + for (i = 0; i < 2; ++i) + ; } #pragma omp master @@ -72,6 +87,9 @@ f_omp (void) ; #pragma acc data /* { dg-error "may not be nested" } */ ; +#pragma acc loop /* { dg-error "may not be closely nested" } */ + for (i = 0; i < 2; ++i) + ; } #pragma omp critical @@ -82,6 +100,9 @@ f_omp (void) ; #pragma acc data /* { dg-error "may not be nested" } */ ; +#pragma acc loop /* { dg-error "may not be closely nested" } */ + for (i = 0; i < 2; ++i) + ; } #pragma omp ordered @@ -92,6 +113,9 @@ f_omp (void) ; #pragma acc data /* { dg-error "may not be nested" } */ ; +#pragma acc loop /* { dg-error "may not be closely nested" } */ + for (i = 0; i < 2; ++i) + ; } } @@ -289,3 +313,77 @@ f_acc_data (void) ; } } + +/* TODO: Some of these should either be allowed or fail with a more sensible + error message. */ +void +f_acc_loop (void) +{ + int i; + +#pragma acc loop + for (i = 0; i < 2; ++i) + { +#pragma omp parallel /* { dg-error "may not be nested" } */ + ; + } + +#pragma acc loop + for (i = 0; i < 2; ++i) + { +#pragma omp for /* { dg-error "may not be nested" } */ + for (i = 0; i < 3; i++) + ; + } + +#pragma acc loop + for (i = 0; i < 2; ++i) + { +#pragma omp sections /* { dg-error "may not be nested" } */ + { + ; + } + } + +#pragma acc loop + for (i = 0; i < 2; ++i) + { +#pragma omp single /* { dg-error "may not be nested" } */ + ; + } + +#pragma acc loop + for (i = 0; i < 2; ++i) + { +#pragma omp task /* { dg-error "may not be nested" } */ + ; + } + +#pragma acc loop + for (i = 0; i < 2; ++i) + { +#pragma omp master /* { dg-error "may not be nested" } */ + ; + } + +#pragma acc loop + for (i = 0; i < 2; ++i) + { +#pragma omp critical /* { dg-error "may not be nested" } */ + ; + } + +#pragma acc loop + for (i = 0; i < 2; ++i) + { +#pragma omp atomic write + i = 0; /* { dg-error "may not be nested" } */ + } + +#pragma acc loop + for (i = 0; i < 2; ++i) + { +#pragma omp ordered /* { dg-error "may not be nested" } */ + ; + } +} diff --git gcc/testsuite/c-c++-common/goacc/clauses-fail.c gcc/testsuite/c-c++-common/goacc/clauses-fail.c index 133bf81..e8e1278 100644 --- gcc/testsuite/c-c++-common/goacc/clauses-fail.c +++ gcc/testsuite/c-c++-common/goacc/clauses-fail.c @@ -1,6 +1,8 @@ void f (void) { + int i; + #pragma acc parallel one /* { dg-error "expected clause before 'one'" } */ ; @@ -9,4 +11,8 @@ f (void) #pragma acc data two /* { dg-error "expected clause before 'two'" } */ ; + +#pragma acc loop deux /* { dg-error "expected clause before 'deux'" } */ + for (i = 0; i < 2; ++i) + ; } diff --git gcc/testsuite/c-c++-common/goacc/nesting-1.c gcc/testsuite/c-c++-common/goacc/nesting-1.c index 3a22292..a489d2d 100644 --- gcc/testsuite/c-c++-common/goacc/nesting-1.c +++ gcc/testsuite/c-c++-common/goacc/nesting-1.c @@ -1,13 +1,56 @@ void +f_acc_parallel (void) +{ +#pragma acc parallel + { + int i; + +#pragma acc loop + for (i = 0; i < 2; ++i) + ; + } +} + + +void +f_acc_kernels (void) +{ +#pragma acc kernels + { + int i; + +#pragma acc loop + for (i = 0; i < 2; ++i) + ; + } +} + + +void f_acc_data (void) { #pragma acc data { + int i; + #pragma acc parallel ; + +#pragma acc parallel + { +#pragma acc loop + for (i = 0; i < 2; ++i) + ; + } + #pragma acc kernels ; + #pragma acc data ; + +#pragma acc loop + for (i = 0; i < 2; ++i) + ; } } diff --git gcc/testsuite/gcc.dg/goacc/sb-1.c gcc/testsuite/gcc.dg/goacc/sb-1.c index 24c88fe..bcb7272 100644 --- gcc/testsuite/gcc.dg/goacc/sb-1.c +++ gcc/testsuite/gcc.dg/goacc/sb-1.c @@ -2,6 +2,8 @@ void foo() { + int l; + bad1: #pragma acc parallel goto bad1; // { dg-error "invalid branch to/from OpenACC structured block" } @@ -9,6 +11,9 @@ void foo() goto bad1; // { dg-error "invalid branch to/from OpenACC structured block" } #pragma acc data goto bad1; // { dg-error "invalid branch to/from OpenACC structured block" } + #pragma acc loop + for (l = 0; l < 2; ++l) + goto bad1; // { dg-error "invalid branch to/from OpenACC structured block" } goto bad2_parallel; // { dg-error "invalid entry to OpenACC structured block" } #pragma acc parallel @@ -28,6 +33,13 @@ void foo() bad2_data: ; } + goto bad2_loop; // { dg-error "invalid entry to OpenACC structured block" } + #pragma acc loop + for (l = 0; l < 2; ++l) + { + bad2_loop: ; + } + #pragma acc parallel { int i; @@ -51,4 +63,13 @@ void foo() for (i = 0; i < 10; ++i) { ok1_data: break; } } + + #pragma acc loop + for (l = 0; l < 2; ++l) + { + int i; + goto ok1_loop; + for (i = 0; i < 10; ++i) + { ok1_loop: break; } + } } diff --git gcc/testsuite/gcc.dg/goacc/sb-3.c gcc/testsuite/gcc.dg/goacc/sb-3.c new file mode 100644 index 0000000..6c2926c --- /dev/null +++ gcc/testsuite/gcc.dg/goacc/sb-3.c @@ -0,0 +1,18 @@ +// { dg-do compile } + +void f (void) +{ + int i, j; +#pragma acc loop + for(i = 1; i < 30; i++) + { + if (i == 7) goto out; // { dg-error "invalid branch to/from OpenACC structured block" } +#pragma acc loop // { dg-error "work-sharing region may not be closely nested inside of work-sharing, critical, ordered, master or explicit task region" } + for(j = 5; j < 10; j++) + { + if (i == 6 && j == 7) goto out; // { dg-error "invalid branch to/from OpenACC structured block" } + } + } + out: + ; +} diff --git gcc/tree-inline.c gcc/tree-inline.c index cdfe35c..5cfda33 100644 --- gcc/tree-inline.c +++ gcc/tree-inline.c @@ -1342,6 +1342,7 @@ remap_gimple_stmt (gimple stmt, copy_body_data *id) break; case GIMPLE_OMP_FOR: + gcc_assert (!is_gimple_omp_oacc_specifically (stmt)); s1 = remap_gimple_seq (gimple_omp_body (stmt), id); s2 = remap_gimple_seq (gimple_omp_for_pre_body (stmt), id); copy = gimple_build_omp_for (s1, gimple_omp_for_kind (stmt), diff --git gcc/tree-nested.c gcc/tree-nested.c index 397f851..e8ba1e3 100644 --- gcc/tree-nested.c +++ gcc/tree-nested.c @@ -622,6 +622,8 @@ walk_gimple_omp_for (gimple for_stmt, walk_stmt_fn callback_stmt, walk_tree_fn callback_op, struct nesting_info *info) { + gcc_assert (!is_gimple_omp_oacc_specifically (for_stmt)); + struct walk_stmt_info wi; gimple_seq seq; tree t; @@ -1282,6 +1284,7 @@ convert_nonlocal_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p, break; case GIMPLE_OMP_FOR: + gcc_assert (!is_gimple_omp_oacc_specifically (stmt)); save_suppress = info->suppress_expansion; convert_nonlocal_omp_clauses (gimple_omp_for_clauses_ptr (stmt), wi); walk_gimple_omp_for (stmt, convert_nonlocal_reference_stmt, @@ -1746,6 +1749,7 @@ convert_local_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p, break; case GIMPLE_OMP_FOR: + gcc_assert (!is_gimple_omp_oacc_specifically (stmt)); save_suppress = info->suppress_expansion; convert_local_omp_clauses (gimple_omp_for_clauses_ptr (stmt), wi); walk_gimple_omp_for (stmt, convert_local_reference_stmt, @@ -2178,6 +2182,7 @@ convert_gimple_call (gimple_stmt_iterator *gsi, bool *handled_ops_p, break; case GIMPLE_OMP_FOR: + gcc_assert (!is_gimple_omp_oacc_specifically (stmt)); walk_body (convert_gimple_call, NULL, info, gimple_omp_for_pre_body_ptr (stmt)); /* FALLTHRU */ diff --git gcc/tree-pretty-print.c gcc/tree-pretty-print.c index 49e5f6c..6c311790 100644 --- gcc/tree-pretty-print.c +++ gcc/tree-pretty-print.c @@ -2538,14 +2538,14 @@ dump_generic_node (pretty_printer *buffer, tree node, int spc, int flags, pp_string (buffer, "#pragma simd"); goto dump_omp_loop; - case OACC_LOOP: - pp_string (buffer, "#pragma acc loop"); - goto dump_omp_loop; - case OMP_DISTRIBUTE: pp_string (buffer, "#pragma omp distribute"); goto dump_omp_loop; + case OACC_LOOP: + pp_string (buffer, "#pragma acc loop"); + goto dump_omp_loop; + case OMP_TEAMS: pp_string (buffer, "#pragma omp teams"); dump_omp_clauses (buffer, OMP_TEAMS_CLAUSES (node), spc, flags); diff --git gcc/tree.def gcc/tree.def index d9e4eb41..a9916f4 100644 --- gcc/tree.def +++ gcc/tree.def @@ -1065,7 +1065,7 @@ DEFTREECODE (OMP_TASK, "omp_task", tcc_statement, 2) private. N1, N2 and INCR are required to be loop invariant integer expressions that are evaluated without any synchronization. The evaluation order, frequency of evaluation and side-effects are - unspecified by the standard. */ + unspecified by the standards. */ DEFTREECODE (OMP_FOR, "omp_for", tcc_statement, 6) /* OpenMP - #pragma omp simd [clause1 ... clauseN] @@ -1076,14 +1076,14 @@ DEFTREECODE (OMP_SIMD, "omp_simd", tcc_statement, 6) Operands like for OMP_FOR. */ DEFTREECODE (CILK_SIMD, "cilk_simd", tcc_statement, 6) -/* OpenACC - #pragma acc loop [clause1 ... clauseN] - Operands like for OMP_FOR. */ -DEFTREECODE (OACC_LOOP, "oacc_loop", tcc_statement, 6) - /* OpenMP - #pragma omp distribute [clause1 ... clauseN] Operands like for OMP_FOR. */ DEFTREECODE (OMP_DISTRIBUTE, "omp_distribute", tcc_statement, 6) +/* OpenMP - #pragma acc loop [clause1 ... clauseN] + Operands like for OMP_FOR. */ +DEFTREECODE (OACC_LOOP, "oacc_loop", tcc_statement, 6) + /* OpenMP - #pragma omp teams [clause1 ... clauseN] Operand 0: OMP_TEAMS_BODY: Teams body. Operand 1: OMP_TEAMS_CLAUSES: List of clauses. */ diff --git gcc/tree.h gcc/tree.h index 6668895..196ec3e 100644 --- gcc/tree.h +++ gcc/tree.h @@ -1210,7 +1210,7 @@ extern void protected_set_expr_location (tree, location_t); #define OMP_TASKREG_BODY(NODE) TREE_OPERAND (OMP_TASKREG_CHECK (NODE), 0) #define OMP_TASKREG_CLAUSES(NODE) TREE_OPERAND (OMP_TASKREG_CHECK (NODE), 1) -#define OMP_LOOP_CHECK(NODE) TREE_RANGE_CHECK (NODE, OMP_FOR, OMP_DISTRIBUTE) +#define OMP_LOOP_CHECK(NODE) TREE_RANGE_CHECK (NODE, OMP_FOR, OACC_LOOP) #define OMP_FOR_BODY(NODE) TREE_OPERAND (OMP_LOOP_CHECK (NODE), 0) #define OMP_FOR_CLAUSES(NODE) TREE_OPERAND (OMP_LOOP_CHECK (NODE), 1) #define OMP_FOR_INIT(NODE) TREE_OPERAND (OMP_LOOP_CHECK (NODE), 2)