From patchwork Thu Jun 11 12:14:20 2015 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Jakub Jelinek X-Patchwork-Id: 483111 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 1C1AD140285 for ; Thu, 11 Jun 2015 22:14:41 +1000 (AEST) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=ItWF/ElG; 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:cc:subject:message-id:reply-to:references:mime-version :content-type:in-reply-to; q=dns; s=default; b=AZjKJI2NIez/sgf88 iAlj8ANnvjozPASFF9ckT21rhk3VLLNzvtpicrpDGr2/x2/Xwifhe2FXgN4YYdsQ ZgA38MUZ5WlBg8CqtLnhNSP02HMzWWPAbCl5lYerj9IPlXjudZHZLzPqclEjgsRa 76F7DSWHFfHMP0Iusppl6i4PMs= 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:cc:subject:message-id:reply-to:references:mime-version :content-type:in-reply-to; s=default; bh=2XNbe3wkTdW3VETds+J0Ehu Y3WA=; b=ItWF/ElGXUH5hwvwIrBmWvGt+rGaEvFY+RZY7G5sO14rKIKzxq2/6yB jlIW+R0zP+UmOFv/CVyXd27njDMj6xLfQS1Y9vZHuNKWGksWNvMzjcPx4nFYjYSN ed5zQdnka8DUqxHxo5MayCQzqMrjzh9XZHE5CS2R3SLDTbyl+zgM= Received: (qmail 23435 invoked by alias); 11 Jun 2015 12:14:31 -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 23424 invoked by uid 89); 11 Jun 2015 12:14:30 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-0.0 required=5.0 tests=AWL, BAYES_50, KAM_LAZY_DOMAIN_SECURITY, SPF_HELO_PASS, T_RP_MATCHES_RCVD autolearn=no 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; Thu, 11 Jun 2015 12:14:28 +0000 Received: from int-mx14.intmail.prod.int.phx2.redhat.com (int-mx14.intmail.prod.int.phx2.redhat.com [10.5.11.27]) by mx1.redhat.com (Postfix) with ESMTPS id 572642D1272; Thu, 11 Jun 2015 12:14:27 +0000 (UTC) Received: from tucnak.zalov.cz (ovpn-116-89.ams2.redhat.com [10.36.116.89]) by int-mx14.intmail.prod.int.phx2.redhat.com (8.14.4/8.14.4) with ESMTP id t5BCEOCw031468 (version=TLSv1/SSLv3 cipher=DHE-RSA-AES256-GCM-SHA384 bits=256 verify=NO); Thu, 11 Jun 2015 08:14:25 -0400 Received: from tucnak.zalov.cz (localhost [127.0.0.1]) by tucnak.zalov.cz (8.14.9/8.14.9) with ESMTP id t5BCEMn6030225; Thu, 11 Jun 2015 14:14:23 +0200 Received: (from jakub@localhost) by tucnak.zalov.cz (8.14.9/8.14.9/Submit) id t5BCELpg030223; Thu, 11 Jun 2015 14:14:21 +0200 Date: Thu, 11 Jun 2015 14:14:20 +0200 From: Jakub Jelinek To: Ilya Verbin , Thomas Schwinge Cc: gcc-patches@gcc.gnu.org, Kirill Yukhin Subject: [gomp4.1] map clause parsing improvements Message-ID: <20150611121420.GY10247@tucnak.redhat.com> Reply-To: Jakub Jelinek References: <20150429111406.GE1751@tucnak.redhat.com> <874mnzrw1z.fsf@schwinge.name> <20150429120644.GG1751@tucnak.redhat.com> <20150609183608.GA47936@msticlxl57.ims.intel.com> MIME-Version: 1.0 Content-Disposition: inline In-Reply-To: <20150609183608.GA47936@msticlxl57.ims.intel.com> User-Agent: Mutt/1.5.23 (2014-03-12) X-IsSubscribed: yes On Tue, Jun 09, 2015 at 09:36:08PM +0300, Ilya Verbin wrote: > On Wed, Apr 29, 2015 at 14:06:44 +0200, Jakub Jelinek wrote: > > [...] The draft requires only alloc or to > > (or always, variants) for enter data and only from or delete (or always, > > variants) for exit data, so in theory it is possible to figure that from > > the call without extra args, but not so for update - enter data is supposed > > to increment reference counts, exit data decrement. [...] > > TR3.pdf also says about 'release' map-type for exit data, but it is not > described in the document. So, I've committed a patch to add parsing release map-kind, and fix up or add verification in C/C++ FE what map-kinds are used. Furthermore, it seems the OpenMP 4.1 always modifier is something completely unrelated to the OpenACC force flag, in OpenMP 4.1 everything is reference count based, and always seems to make a difference only for from/to/tofrom, where it says that the copying is done unconditionally; thus the patch uses a different bit for that. For array sections resulting in GOMP_MAP_POINTER kind perhaps we'll also need an always variant for it, and supposedly on the exit data construct we'll want to turn GOMP_MAP_POINTER into GOMP_MAP_RELEASE or GOMP_MAP_DELETE depending on what it was originally. 2015-06-11 Jakub Jelinek include/ * gomp-constants.h (GOMP_MAP_FLAG_ALWAYS): Define. (enum gomp_map_kind): Add GOMP_MAP_ALWAYS_TO, GOMP_MAP_ALWAYS_FROM, GOMP_MAP_ALWAYS_TOFROM, GOMP_MAP_DELETE, GOMP_MAP_RELEASE. gcc/ * omp-low.c (lower_omp_target): Accept GOMP_MAP_RELEASE, GOMP_MAP_ALWAYS_TO, GOMP_MAP_ALWAYS_FROM and GOMP_MAP_ALWAYS_TOFROM. Accept GOMP_MAP_FORCE* except for FORCE_DEALLOC only for OpenACC. * tree-pretty-print.c (dump_omp_clause): Print GOMP_MAP_FORCE_* as force_*, handle GOMP_MAP_ALWAYS_*. gcc/c/ * c-parser.c (c_parser_omp_clause_map): Handle release map kind. Adjust to use GOMP_MAP_ALWAYS_* and only on clauses where it makes sense. (c_parser_omp_target_data): Diagnose invalid map-kind for the construct. (c_parser_omp_target_enter_data): Adjust for new encoding of always,from, always,to and always,tofrom. Accept GOMP_MAP_POINTER. (c_parser_omp_target_exit_data): Likewise. (c_parser_omp_target): Diagnose invalid map-kind for the construct. gcc/cp/ * parser.c (cp_parser_omp_clause_map): Handle release map kind. Adjust to use GOMP_MAP_ALWAYS_* and only on clauses where it makes sense. (cp_parser_omp_target_data): Diagnose invalid map-kind for the construct. (cp_parser_omp_target_enter_data): Adjust for new encoding of always,from, always,to and always,tofrom. Accept GOMP_MAP_POINTER. (cp_parser_omp_target_exit_data): Likewise. (cp_parser_omp_target): Diagnose invalid map-kind for the construct. Jakub --- include/gomp-constants.h.jj 2015-05-21 11:12:09.000000000 +0200 +++ include/gomp-constants.h 2015-06-11 11:24:32.041654947 +0200 @@ -41,6 +41,8 @@ #define GOMP_MAP_FLAG_SPECIAL_1 (1 << 3) #define GOMP_MAP_FLAG_SPECIAL (GOMP_MAP_FLAG_SPECIAL_1 \ | GOMP_MAP_FLAG_SPECIAL_0) +/* OpenMP always flag. */ +#define GOMP_MAP_FLAG_ALWAYS (1 << 6) /* Flag to force a specific behavior (or else, trigger a run-time error). */ #define GOMP_MAP_FLAG_FORCE (1 << 7) @@ -77,7 +79,21 @@ enum gomp_map_kind /* ..., and copy from device. */ GOMP_MAP_FORCE_FROM = (GOMP_MAP_FLAG_FORCE | GOMP_MAP_FROM), /* ..., and copy to and from device. */ - GOMP_MAP_FORCE_TOFROM = (GOMP_MAP_FLAG_FORCE | GOMP_MAP_TOFROM) + GOMP_MAP_FORCE_TOFROM = (GOMP_MAP_FLAG_FORCE | GOMP_MAP_TOFROM), + /* If not already present, allocate. And unconditionally copy to + device. */ + GOMP_MAP_ALWAYS_TO = (GOMP_MAP_FLAG_ALWAYS | GOMP_MAP_TO), + /* If not already present, allocate. And unconditionally copy from + device. */ + GOMP_MAP_ALWAYS_FROM = (GOMP_MAP_FLAG_ALWAYS | GOMP_MAP_FROM), + /* If not already present, allocate. And unconditionally copy to and from + device. */ + GOMP_MAP_ALWAYS_TOFROM = (GOMP_MAP_FLAG_ALWAYS | GOMP_MAP_TOFROM), + /* OpenMP 4.1 alias for forced deallocation. */ + GOMP_MAP_DELETE = GOMP_MAP_FORCE_DEALLOC, + /* Decrement usage count and deallocate if zero. */ + GOMP_MAP_RELEASE = (GOMP_MAP_FLAG_ALWAYS + | GOMP_MAP_FORCE_DEALLOC) }; #define GOMP_MAP_COPY_TO_P(X) \ --- gcc/omp-low.c.jj 2015-06-10 19:50:26.000000000 +0200 +++ gcc/omp-low.c 2015-06-11 11:35:02.515892363 +0200 @@ -12511,13 +12511,17 @@ lower_omp_target (gimple_stmt_iterator * case GOMP_MAP_TOFROM: case GOMP_MAP_POINTER: case GOMP_MAP_TO_PSET: + case GOMP_MAP_FORCE_DEALLOC: + case GOMP_MAP_RELEASE: + case GOMP_MAP_ALWAYS_TO: + case GOMP_MAP_ALWAYS_FROM: + case GOMP_MAP_ALWAYS_TOFROM: + break; case GOMP_MAP_FORCE_ALLOC: case GOMP_MAP_FORCE_TO: case GOMP_MAP_FORCE_FROM: case GOMP_MAP_FORCE_TOFROM: case GOMP_MAP_FORCE_PRESENT: - case GOMP_MAP_FORCE_DEALLOC: - break; case GOMP_MAP_FORCE_DEVICEPTR: gcc_assert (is_gimple_omp_oacc (stmt)); break; --- gcc/tree-pretty-print.c.jj 2015-05-19 18:56:43.000000000 +0200 +++ gcc/tree-pretty-print.c 2015-06-11 11:32:39.814102121 +0200 @@ -560,26 +560,38 @@ dump_omp_clause (pretty_printer *pp, tre pp_string (pp, "tofrom"); break; case GOMP_MAP_FORCE_ALLOC: - pp_string (pp, "always,alloc"); + pp_string (pp, "force_alloc"); break; case GOMP_MAP_FORCE_TO: - pp_string (pp, "always,to"); + pp_string (pp, "force_to"); break; case GOMP_MAP_FORCE_FROM: - pp_string (pp, "always,from"); + pp_string (pp, "force_from"); break; case GOMP_MAP_FORCE_TOFROM: - pp_string (pp, "always,tofrom"); + pp_string (pp, "force_tofrom"); break; case GOMP_MAP_FORCE_PRESENT: pp_string (pp, "force_present"); break; case GOMP_MAP_FORCE_DEALLOC: - pp_string (pp, "always,delete"); + pp_string (pp, "delete"); break; case GOMP_MAP_FORCE_DEVICEPTR: pp_string (pp, "force_deviceptr"); break; + case GOMP_MAP_ALWAYS_TO: + pp_string (pp, "always,to"); + break; + case GOMP_MAP_ALWAYS_FROM: + pp_string (pp, "always,from"); + break; + case GOMP_MAP_ALWAYS_TOFROM: + pp_string (pp, "always,tofrom"); + break; + case GOMP_MAP_RELEASE: + pp_string (pp, "release"); + break; default: gcc_unreachable (); } --- gcc/c/c-parser.c.jj 2015-06-10 14:52:06.000000000 +0200 +++ gcc/c/c-parser.c 2015-06-11 13:02:30.788629315 +0200 @@ -11661,7 +11661,7 @@ c_parser_omp_clause_depend (c_parser *pa OpenMP 4.1: map-kind: - alloc | to | from | tofrom | delete + alloc | to | from | tofrom | release | delete map ( always [,] map-kind: variable-list ) */ @@ -11702,6 +11702,7 @@ c_parser_omp_clause_map (c_parser *parse || strcmp ("to", p) == 0 || strcmp ("from", p) == 0 || strcmp ("tofrom", p) == 0 + || strcmp ("release", p) == 0 || strcmp ("delete", p) == 0) { c_parser_consume_token (parser); @@ -11718,13 +11719,15 @@ c_parser_omp_clause_map (c_parser *parse if (strcmp ("alloc", p) == 0) kind = GOMP_MAP_ALLOC; else if (strcmp ("to", p) == 0) - kind = GOMP_MAP_TO; + kind = always ? GOMP_MAP_ALWAYS_TO : GOMP_MAP_TO; else if (strcmp ("from", p) == 0) - kind = GOMP_MAP_FROM; + kind = always ? GOMP_MAP_ALWAYS_FROM : GOMP_MAP_FROM; else if (strcmp ("tofrom", p) == 0) - kind = GOMP_MAP_TOFROM; + kind = always ? GOMP_MAP_ALWAYS_TOFROM : GOMP_MAP_TOFROM; + else if (strcmp ("release", p) == 0) + kind = GOMP_MAP_RELEASE; else if (strcmp ("delete", p) == 0) - kind = GOMP_MAP_FORCE_DEALLOC; + kind = GOMP_MAP_DELETE; else { c_parser_error (parser, "invalid map kind"); @@ -11732,8 +11735,6 @@ c_parser_omp_clause_map (c_parser *parse "expected %<)%>"); return list; } - if (always) - kind = (enum gomp_map_kind) (kind | GOMP_MAP_FLAG_FORCE); c_parser_consume_token (parser); c_parser_consume_token (parser); } @@ -14237,11 +14238,40 @@ c_parser_omp_target_data (location_t loc tree clauses = c_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK, "#pragma omp target data"); - if (find_omp_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE) + int map_seen = 0; + for (tree *pc = &clauses; *pc;) { - error_at (loc, - "%<#pragma omp target data%> must contain at least one " - "% clause"); + if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP) + switch (OMP_CLAUSE_MAP_KIND (*pc)) + { + case GOMP_MAP_TO: + case GOMP_MAP_ALWAYS_TO: + case GOMP_MAP_FROM: + case GOMP_MAP_ALWAYS_FROM: + case GOMP_MAP_TOFROM: + case GOMP_MAP_ALWAYS_TOFROM: + case GOMP_MAP_ALLOC: + case GOMP_MAP_POINTER: + map_seen = 3; + break; + default: + map_seen |= 1; + error_at (OMP_CLAUSE_LOCATION (*pc), + "%<#pragma omp target data%> with map-type other " + "than %, %, % or % " + "on % clause"); + *pc = OMP_CLAUSE_CHAIN (*pc); + continue; + } + pc = &OMP_CLAUSE_CHAIN (*pc); + } + + if (map_seen != 3) + { + if (map_seen == 0) + error_at (loc, + "%<#pragma omp target data%> must contain at least " + "one % clause"); return NULL_TREE; } @@ -14348,10 +14378,12 @@ c_parser_omp_target_enter_data (location for (tree *pc = &clauses; *pc;) { if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP) - switch (OMP_CLAUSE_MAP_KIND (*pc) & ~GOMP_MAP_FLAG_FORCE) + switch (OMP_CLAUSE_MAP_KIND (*pc)) { case GOMP_MAP_TO: + case GOMP_MAP_ALWAYS_TO: case GOMP_MAP_ALLOC: + case GOMP_MAP_POINTER: map_seen = 3; break; default: @@ -14430,17 +14462,21 @@ c_parser_omp_target_exit_data (location_ for (tree *pc = &clauses; *pc;) { if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP) - switch (OMP_CLAUSE_MAP_KIND (*pc) & ~GOMP_MAP_FLAG_FORCE) + switch (OMP_CLAUSE_MAP_KIND (*pc)) { case GOMP_MAP_FROM: - case GOMP_MAP_FORCE_DEALLOC & ~GOMP_MAP_FLAG_FORCE: + case GOMP_MAP_ALWAYS_FROM: + case GOMP_MAP_RELEASE: + case GOMP_MAP_DELETE: + case GOMP_MAP_POINTER: map_seen = 3; break; default: map_seen |= 1; error_at (OMP_CLAUSE_LOCATION (*pc), "%<#pragma omp target exit data%> with map-type other " - "than % or % on % clause"); + "than %, % or % on %" + " clause"); *pc = OMP_CLAUSE_CHAIN (*pc); continue; } @@ -14480,6 +14516,7 @@ c_parser_omp_target (c_parser *parser, e { location_t loc = c_parser_peek_token (parser)->location; c_parser_consume_pragma (parser); + tree *pc = NULL, stmt, block; if (context != pragma_stmt && context != pragma_compound) { @@ -14519,7 +14556,8 @@ c_parser_omp_target (c_parser *parser, e OMP_TARGET_CLAUSES (stmt) = cclauses[C_OMP_CLAUSE_SPLIT_TARGET]; OMP_TARGET_BODY (stmt) = block; add_stmt (stmt); - return true; + pc = &OMP_TARGET_CLAUSES (stmt); + goto check_clauses; } else if (!flag_openmp) /* flag_openmp_simd */ { @@ -14551,19 +14589,46 @@ c_parser_omp_target (c_parser *parser, e } } - tree stmt = make_node (OMP_TARGET); + stmt = make_node (OMP_TARGET); TREE_TYPE (stmt) = void_type_node; OMP_TARGET_CLAUSES (stmt) = c_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK, "#pragma omp target"); + pc = &OMP_TARGET_CLAUSES (stmt); keep_next_level (); - tree block = c_begin_compound_stmt (true); + block = c_begin_compound_stmt (true); add_stmt (c_parser_omp_structured_block (parser)); OMP_TARGET_BODY (stmt) = c_end_compound_stmt (loc, block, true); SET_EXPR_LOCATION (stmt, loc); add_stmt (stmt); + +check_clauses: + while (*pc) + { + if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP) + switch (OMP_CLAUSE_MAP_KIND (*pc)) + { + case GOMP_MAP_TO: + case GOMP_MAP_ALWAYS_TO: + case GOMP_MAP_FROM: + case GOMP_MAP_ALWAYS_FROM: + case GOMP_MAP_TOFROM: + case GOMP_MAP_ALWAYS_TOFROM: + case GOMP_MAP_ALLOC: + case GOMP_MAP_POINTER: + break; + default: + error_at (OMP_CLAUSE_LOCATION (*pc), + "%<#pragma omp target%> with map-type other " + "than %, %, % or % " + "on % clause"); + *pc = OMP_CLAUSE_CHAIN (*pc); + continue; + } + pc = &OMP_CLAUSE_CHAIN (*pc); + } return true; } --- gcc/cp/parser.c.jj 2015-06-03 19:48:09.000000000 +0200 +++ gcc/cp/parser.c 2015-06-11 13:03:19.746878152 +0200 @@ -29150,7 +29150,7 @@ cp_parser_omp_clause_depend (cp_parser * OpenMP 4.1: map-kind: - alloc | to | from | tofrom | delete + alloc | to | from | tofrom | release | delete map ( always [,] map-kind: variable-list ) */ @@ -29195,13 +29195,15 @@ cp_parser_omp_clause_map (cp_parser *par if (strcmp ("alloc", p) == 0) kind = GOMP_MAP_ALLOC; else if (strcmp ("to", p) == 0) - kind = GOMP_MAP_TO; + kind = always ? GOMP_MAP_ALWAYS_TO : GOMP_MAP_TO; else if (strcmp ("from", p) == 0) - kind = GOMP_MAP_FROM; + kind = always ? GOMP_MAP_ALWAYS_FROM : GOMP_MAP_FROM; else if (strcmp ("tofrom", p) == 0) - kind = GOMP_MAP_TOFROM; + kind = always ? GOMP_MAP_ALWAYS_TOFROM : GOMP_MAP_TOFROM; + else if (strcmp ("release", p) == 0) + kind = GOMP_MAP_RELEASE; else if (strcmp ("delete", p) == 0) - kind = GOMP_MAP_FORCE_DEALLOC; + kind = GOMP_MAP_DELETE; else { cp_parser_error (parser, "invalid map kind"); @@ -29210,8 +29212,6 @@ cp_parser_omp_clause_map (cp_parser *par /*consume_paren=*/true); return list; } - if (always) - kind = (enum gomp_map_kind) (kind | GOMP_MAP_FLAG_FORCE); cp_lexer_consume_token (parser->lexer); cp_lexer_consume_token (parser->lexer); } @@ -31690,11 +31690,40 @@ cp_parser_omp_target_data (cp_parser *pa tree clauses = cp_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK, "#pragma omp target data", pragma_tok); - if (find_omp_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE) + int map_seen = 0; + for (tree *pc = &clauses; *pc;) { - error_at (pragma_tok->location, - "%<#pragma omp target data%> must contain at least one " - "% clause"); + if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP) + switch (OMP_CLAUSE_MAP_KIND (*pc)) + { + case GOMP_MAP_TO: + case GOMP_MAP_ALWAYS_TO: + case GOMP_MAP_FROM: + case GOMP_MAP_ALWAYS_FROM: + case GOMP_MAP_TOFROM: + case GOMP_MAP_ALWAYS_TOFROM: + case GOMP_MAP_ALLOC: + case GOMP_MAP_POINTER: + map_seen = 3; + break; + default: + map_seen |= 1; + error_at (OMP_CLAUSE_LOCATION (*pc), + "%<#pragma omp target data%> with map-type other " + "than %, %, % or % " + "on % clause"); + *pc = OMP_CLAUSE_CHAIN (*pc); + continue; + } + pc = &OMP_CLAUSE_CHAIN (*pc); + } + + if (map_seen != 3) + { + if (map_seen == 0) + error_at (pragma_tok->location, + "%<#pragma omp target data%> must contain at least " + "one % clause"); return NULL_TREE; } @@ -31759,10 +31788,12 @@ cp_parser_omp_target_enter_data (cp_pars for (tree *pc = &clauses; *pc;) { if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP) - switch (OMP_CLAUSE_MAP_KIND (*pc) & ~GOMP_MAP_FLAG_FORCE) + switch (OMP_CLAUSE_MAP_KIND (*pc)) { case GOMP_MAP_TO: + case GOMP_MAP_ALWAYS_TO: case GOMP_MAP_ALLOC: + case GOMP_MAP_POINTER: map_seen = 3; break; default: @@ -31842,17 +31873,21 @@ cp_parser_omp_target_exit_data (cp_parse for (tree *pc = &clauses; *pc;) { if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP) - switch (OMP_CLAUSE_MAP_KIND (*pc) & ~GOMP_MAP_FLAG_FORCE) + switch (OMP_CLAUSE_MAP_KIND (*pc)) { case GOMP_MAP_FROM: - case GOMP_MAP_FORCE_DEALLOC & ~GOMP_MAP_FLAG_FORCE: + case GOMP_MAP_ALWAYS_FROM: + case GOMP_MAP_RELEASE: + case GOMP_MAP_DELETE: + case GOMP_MAP_POINTER: map_seen = 3; break; default: map_seen |= 1; error_at (OMP_CLAUSE_LOCATION (*pc), "%<#pragma omp target exit data%> with map-type other " - "than % or % on % clause"); + "than %, % or % on %" + " clause"); *pc = OMP_CLAUSE_CHAIN (*pc); continue; } @@ -31934,6 +31969,8 @@ static bool cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, enum pragma_context context) { + tree *pc = NULL, stmt; + if (context != pragma_stmt && context != pragma_compound) { cp_parser_error (parser, "expected declaration specifiers"); @@ -31975,7 +32012,8 @@ cp_parser_omp_target (cp_parser *parser, OMP_TARGET_CLAUSES (stmt) = cclauses[C_OMP_CLAUSE_SPLIT_TARGET]; OMP_TARGET_BODY (stmt) = body; add_stmt (stmt); - return true; + pc = &OMP_TARGET_CLAUSES (stmt); + goto check_clauses; } else if (!flag_openmp) /* flag_openmp_simd */ { @@ -32007,17 +32045,44 @@ cp_parser_omp_target (cp_parser *parser, } } - tree stmt = make_node (OMP_TARGET); + stmt = make_node (OMP_TARGET); TREE_TYPE (stmt) = void_type_node; OMP_TARGET_CLAUSES (stmt) = cp_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK, "#pragma omp target", pragma_tok); + pc = &OMP_TARGET_CLAUSES (stmt); keep_next_level (true); OMP_TARGET_BODY (stmt) = cp_parser_omp_structured_block (parser); SET_EXPR_LOCATION (stmt, pragma_tok->location); add_stmt (stmt); + +check_clauses: + while (*pc) + { + if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP) + switch (OMP_CLAUSE_MAP_KIND (*pc)) + { + case GOMP_MAP_TO: + case GOMP_MAP_ALWAYS_TO: + case GOMP_MAP_FROM: + case GOMP_MAP_ALWAYS_FROM: + case GOMP_MAP_TOFROM: + case GOMP_MAP_ALWAYS_TOFROM: + case GOMP_MAP_ALLOC: + case GOMP_MAP_POINTER: + break; + default: + error_at (OMP_CLAUSE_LOCATION (*pc), + "%<#pragma omp target%> with map-type other " + "than %, %, % or % " + "on % clause"); + *pc = OMP_CLAUSE_CHAIN (*pc); + continue; + } + pc = &OMP_CLAUSE_CHAIN (*pc); + } return true; }