From patchwork Wed Nov 5 16:36:46 2014 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 407083 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 6862914003E for ; Thu, 6 Nov 2014 03:37:18 +1100 (AEDT) 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=jvuFWLMAeG08k0II iBtWKzO2IvFNeyPvwpjo/sELMgwEaw8+0Ex6vF0fgfcUbWXGZsudb8+Obj0AjYj8 iYkPL+6E6o3TukLKwynZBc+RahmQVTff76S9JrWKvlYcxYt3YOVYBheVatQ0U+Wa N0WeeXKovAwEi2K6WHSe8MN3a9E= 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=dNokTPV68Y4R6o+vYfkjXl lL69Y=; b=nZCv2HU+btlqa05zxQUuPSwEiMgRJI5rNmAHVMRIQl1Y/1BHbiPeRQ AQ3+AWvibRlquicD/C43IrhunaAP7kD5cOZT4xZgjYfuA1wgJW/rF3QvZpt6U8Hw NAMwVUq/CK0bn3HUsnZYPB73xhI5ZyQPxh5Mi5N9u6l07gU0uUBzk= Received: (qmail 13214 invoked by alias); 5 Nov 2014 16:36:59 -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 13014 invoked by uid 89); 5 Nov 2014 16:36:57 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.0 required=5.0 tests=AWL, BAYES_00, RCVD_IN_DNSWL_NONE 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; Wed, 05 Nov 2014 16:36:54 +0000 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-FEM-01.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1Xm3Zi-0004uG-Hc from Thomas_Schwinge@mentor.com for gcc-patches@gcc.gnu.org; Wed, 05 Nov 2014 08:36:51 -0800 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.3.181.6; Wed, 5 Nov 2014 16:36:49 +0000 From: Thomas Schwinge To: CC: James Norris , Cesar Philippidis Subject: [gomp4] OpenACC cache directive maintenance (was: [PATCH 4/6] [GOMP4] OpenACC 1.0+ support in fortran front-end) In-Reply-To: <874n4tqink.fsf@schwinge.name> References: <52E158EF.9050009@samsung.com> <52E1595D.9000007@samsung.com> <52E1597E.8070809@samsung.com> <52E15999.6010505@samsung.com> <52E159BD.9040407@samsung.com> <874n4tqink.fsf@schwinge.name> <87d291ehls.fsf@kepler.schwinge.homeip.net> User-Agent: Notmuch/0.9-101-g81dad07 (http://notmuchmail.org) Emacs/24.3.1 (i586-pc-linux-gnu) Date: Wed, 5 Nov 2014 17:36:46 +0100 Message-ID: <87a945eh9d.fsf@kepler.schwinge.homeip.net> MIME-Version: 1.0 Hi! On Fri, 24 Jan 2014 20:33:35 +0100, I wrote: > On Thu, 23 Jan 2014 22:04:45 +0400, Ilmir Usmanov wrote: > > Subject: [PATCH 4/6] OpenACC GENERIC nodes > > --- a/gcc/tree-core.h > > +++ b/gcc/tree-core.h > > @@ -216,12 +216,18 @@ enum omp_clause_code { > > + /* Internal structure to hold OpenACC cache directive's variable-list. > > + #pragma acc cache (variable-_ist). */ > > + OACC_NO_CLAUSE_CACHE, > > Hmm, yeah, while *_NO_CLAUSE_* perhaps isn't the most beautiful approach, > I think it's fine at least for now. In r217146, I applied the following to gomp-4_0-branch: commit e8e44b733808997d06c0cdf9bf5756ce03530f42 Author: tschwinge Date: Wed Nov 5 16:35:30 2014 +0000 OpenACC cache directive maintenance. gcc/c/ * c-parser.c (c_parser_oacc_cache): Generate OACC_CACHE. * c-typeck.c (c_finish_omp_clauses): Handle OMP_CLAUSE__CACHE_. gcc/cp/ * parser.c (cp_parser_oacc_cache): Generate OACC_CACHE. * semantics.c (finish_omp_clauses): Handle OMP_CLAUSE__CACHE_. gcc/ * gimplify.c (gimplify_oacc_cache): New function. (gimplify_expr): Use it for OACC_CACHE. (gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses): Handle OMP_CLAUSE__CACHE_. gcc/c/ * c-parser.c (c_parser_omp_variable_list) : Remove explicit mark_exp_read invocations. gcc/cp/ * parser.c (cp_parser_omp_var_list_no_open) : Remove explicit mark_exp_read invocations. gcc/ * tree-core.h (enum omp_clause_code): Move OMP_NO_CLAUSE_CACHE next to, and handle it like a data clause. Rename it to OMP_CLAUSE__CACHE_. Update all users. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@217146 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/ChangeLog.gomp | 9 +++++++++ gcc/c/ChangeLog.gomp | 8 ++++++++ gcc/c/c-parser.c | 23 +++++++++++++++-------- gcc/c/c-typeck.c | 1 + gcc/cp/ChangeLog.gomp | 6 ++++++ gcc/cp/parser.c | 24 +++++++++++++++--------- gcc/cp/semantics.c | 1 + gcc/fortran/trans-openmp.c | 2 +- gcc/gimplify.c | 25 ++++++++++++++++++++++--- gcc/omp-low.c | 4 ++-- gcc/tree-core.h | 8 ++++---- gcc/tree-pretty-print.c | 11 +++++++---- gcc/tree.c | 6 +++--- gcc/tree.def | 5 +++-- gcc/tree.h | 2 +- 15 files changed, 98 insertions(+), 37 deletions(-) Grüße, Thomas diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp index fc624c8..2c2b349 100644 --- gcc/ChangeLog.gomp +++ gcc/ChangeLog.gomp @@ -1,5 +1,14 @@ 2014-11-05 Thomas Schwinge + * gimplify.c (gimplify_oacc_cache): New function. + (gimplify_expr): Use it for OACC_CACHE. + (gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses): Handle + OMP_CLAUSE__CACHE_. + + * tree-core.h (enum omp_clause_code): Move OMP_NO_CLAUSE_CACHE + next to, and handle it like a data clause. Rename it to + OMP_CLAUSE__CACHE_. Update all users. + * invoke.texi: Update for OpenACC. * sourcebuild.texi: Likewise. diff --git gcc/c/ChangeLog.gomp gcc/c/ChangeLog.gomp index 7acd7b3..70278b9 100644 --- gcc/c/ChangeLog.gomp +++ gcc/c/ChangeLog.gomp @@ -1,3 +1,11 @@ +2014-11-05 Thomas Schwinge + + * c-parser.c (c_parser_oacc_cache): Generate OACC_CACHE. + * c-typeck.c (c_finish_omp_clauses): Handle OMP_CLAUSE__CACHE_. + + * c-parser.c (c_parser_omp_variable_list) : + Remove explicit mark_exp_read invocations. + 2014-11-05 James Norris * c-parser.c (c_parser_omp_variable_list): Handle diff --git gcc/c/c-parser.c gcc/c/c-parser.c index 410b19f..40d4314 100644 --- gcc/c/c-parser.c +++ gcc/c/c-parser.c @@ -10053,7 +10053,7 @@ c_parser_omp_variable_list (c_parser *parser, { switch (kind) { - case OMP_NO_CLAUSE_CACHE: + case OMP_CLAUSE__CACHE_: if (c_parser_peek_token (parser)->type != CPP_OPEN_SQUARE) { c_parser_error (parser, "expected %<[%>"); @@ -10100,11 +10100,8 @@ c_parser_omp_variable_list (c_parser *parser, break; } - if (kind == OMP_NO_CLAUSE_CACHE) + if (kind == OMP_CLAUSE__CACHE_) { - mark_exp_read (low_bound); - mark_exp_read (length); - if (TREE_CODE (low_bound) != INTEGER_CST && !TREE_READONLY (low_bound)) { @@ -11901,12 +11898,22 @@ c_parser_omp_structured_block (c_parser *parser) */ static tree -c_parser_oacc_cache (location_t loc __attribute__((unused)), c_parser *parser) +c_parser_oacc_cache (location_t loc, c_parser *parser) { - c_parser_omp_var_list_parens (parser, OMP_NO_CLAUSE_CACHE, NULL); + tree stmt, clauses; + + clauses = c_parser_omp_var_list_parens (parser, OMP_CLAUSE__CACHE_, NULL); + clauses = c_finish_omp_clauses (clauses); + c_parser_skip_to_pragma_eol (parser); - return NULL_TREE; + stmt = make_node (OACC_CACHE); + TREE_TYPE (stmt) = void_type_node; + OACC_CACHE_CLAUSES (stmt) = clauses; + SET_EXPR_LOCATION (stmt, loc); + add_stmt (stmt); + + return stmt; } /* OpenACC 2.0: diff --git gcc/c/c-typeck.c gcc/c/c-typeck.c index 76503e4..e315690d 100644 --- gcc/c/c-typeck.c +++ gcc/c/c-typeck.c @@ -12204,6 +12204,7 @@ c_finish_omp_clauses (tree clauses) case OMP_CLAUSE_MAP: case OMP_CLAUSE_TO: case OMP_CLAUSE_FROM: + case OMP_CLAUSE__CACHE_: t = OMP_CLAUSE_DECL (c); if (TREE_CODE (t) == TREE_LIST) { diff --git gcc/cp/ChangeLog.gomp gcc/cp/ChangeLog.gomp index 024e6a5..46d4912 100644 --- gcc/cp/ChangeLog.gomp +++ gcc/cp/ChangeLog.gomp @@ -1,5 +1,11 @@ 2014-11-05 Thomas Schwinge + * parser.c (cp_parser_oacc_cache): Generate OACC_CACHE. + * semantics.c (finish_omp_clauses): Handle OMP_CLAUSE__CACHE_. + + * parser.c (cp_parser_omp_var_list_no_open) : + Remove explicit mark_exp_read invocations. + * parser.c (cp_parser_omp_clause_name): Also look for "pcopy", "pcopyin", "pcopyout", "pcreate". Look for "wait" instead of "WAIT". diff --git gcc/cp/parser.c gcc/cp/parser.c index 3ef2de7..ea4ad2f 100644 --- gcc/cp/parser.c +++ gcc/cp/parser.c @@ -27669,7 +27669,7 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind, { switch (kind) { - case OMP_NO_CLAUSE_CACHE: + case OMP_CLAUSE__CACHE_: if (cp_lexer_peek_token (parser->lexer)->type != CPP_OPEN_SQUARE) { error_at (token->location, "expected %<[%>"); @@ -27708,11 +27708,8 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind, RT_CLOSE_SQUARE)) goto skip_comma; - if (kind == OMP_NO_CLAUSE_CACHE) + if (kind == OMP_CLAUSE__CACHE_) { - mark_exp_read (low_bound); - mark_exp_read (length); - if (TREE_CODE (low_bound) != INTEGER_CST && !TREE_READONLY (low_bound)) { @@ -31410,13 +31407,22 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, */ static tree -cp_parser_oacc_cache (cp_parser *parser, - cp_token *pragma_tok __attribute__((unused))) +cp_parser_oacc_cache (cp_parser *parser, cp_token *pragma_tok) { - cp_parser_omp_var_list (parser, OMP_NO_CLAUSE_CACHE, NULL_TREE); + tree stmt, clauses; + + clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE__CACHE_, NULL_TREE); + clauses = finish_omp_clauses (clauses); + cp_parser_require_pragma_eol (parser, cp_lexer_peek_token (parser->lexer)); - return NULL_TREE; + stmt = make_node (OACC_CACHE); + TREE_TYPE (stmt) = void_type_node; + OACC_CACHE_CLAUSES (stmt) = clauses; + SET_EXPR_LOCATION (stmt, pragma_tok->location); + add_stmt (stmt); + + return stmt; } /* OpenACC 2.0: diff --git gcc/cp/semantics.c gcc/cp/semantics.c index 2457a6f..6e35eef 100644 --- gcc/cp/semantics.c +++ gcc/cp/semantics.c @@ -5704,6 +5704,7 @@ finish_omp_clauses (tree clauses) case OMP_CLAUSE_MAP: case OMP_CLAUSE_TO: case OMP_CLAUSE_FROM: + case OMP_CLAUSE__CACHE_: t = OMP_CLAUSE_DECL (c); if (TREE_CODE (t) == TREE_LIST) { diff --git gcc/fortran/trans-openmp.c gcc/fortran/trans-openmp.c index 97613ae..7dd4498 100644 --- gcc/fortran/trans-openmp.c +++ gcc/fortran/trans-openmp.c @@ -1807,7 +1807,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, clause_code = OMP_CLAUSE_DEVICE_RESIDENT; goto add_clause; case OMP_LIST_CACHE: - clause_code = OMP_NO_CLAUSE_CACHE; + clause_code = OMP_CLAUSE__CACHE_; goto add_clause; add_clause: diff --git gcc/gimplify.c gcc/gimplify.c index bfd7f66..d58876f 100644 --- gcc/gimplify.c +++ gcc/gimplify.c @@ -6114,6 +6114,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, case OMP_CLAUSE_TO: case OMP_CLAUSE_FROM: + case OMP_CLAUSE__CACHE_: decl = OMP_CLAUSE_DECL (c); if (error_operand_p (decl)) { @@ -6294,7 +6295,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, case OMP_CLAUSE_GANG: case OMP_CLAUSE_ASYNC: case OMP_CLAUSE_WAIT: - case OMP_NO_CLAUSE_CACHE: case OMP_CLAUSE_INDEPENDENT: case OMP_CLAUSE_WORKER: case OMP_CLAUSE_VECTOR: @@ -6641,6 +6641,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, tree *list_p) case OMP_CLAUSE_TO: case OMP_CLAUSE_FROM: + case OMP_CLAUSE__CACHE_: decl = OMP_CLAUSE_DECL (c); if (!DECL_P (decl)) break; @@ -6698,7 +6699,6 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, tree *list_p) case OMP_CLAUSE_GANG: case OMP_CLAUSE_ASYNC: case OMP_CLAUSE_WAIT: - case OMP_NO_CLAUSE_CACHE: case OMP_CLAUSE_INDEPENDENT: case OMP_CLAUSE_WORKER: case OMP_CLAUSE_VECTOR: @@ -6722,6 +6722,21 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, tree *list_p) delete_omp_context (ctx); } +/* Gimplify OACC_CACHE. */ + +static void +gimplify_oacc_cache (tree *expr_p, gimple_seq *pre_p) +{ + tree expr = *expr_p; + + gimplify_scan_omp_clauses (&OACC_CACHE_CLAUSES (expr), pre_p, ORT_WORKSHARE); + gimplify_adjust_omp_clauses (pre_p, &OACC_CACHE_CLAUSES (expr)); + + /* TODO: Do something sensible with this information. */ + + *expr_p = NULL_TREE; +} + /* Gimplify the contents of an OMP_PARALLEL statement. This involves gimplification of the body, as well as scanning the body for used variables. We need to do this scan now, because variable-sized @@ -8312,7 +8327,6 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p, case OACC_HOST_DATA: case OACC_DECLARE: - case OACC_CACHE: sorry ("directive not yet implemented"); ret = GS_ALL_DONE; break; @@ -8352,6 +8366,11 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p, ret = GS_ALL_DONE; break; + case OACC_CACHE: + gimplify_oacc_cache (expr_p, pre_p); + ret = GS_ALL_DONE; + break; + case OACC_DATA: case OMP_SECTIONS: case OMP_SINGLE: diff --git gcc/omp-low.c gcc/omp-low.c index 49cf1ab..1c9d942 100644 --- gcc/omp-low.c +++ gcc/omp-low.c @@ -1982,7 +1982,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_DEVICE_RESIDENT: case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_GANG: - case OMP_NO_CLAUSE_CACHE: + case OMP_CLAUSE__CACHE_: case OMP_CLAUSE_INDEPENDENT: case OMP_CLAUSE_WORKER: case OMP_CLAUSE_VECTOR: @@ -2130,7 +2130,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_DEVICE_RESIDENT: case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_GANG: - case OMP_NO_CLAUSE_CACHE: + case OMP_CLAUSE__CACHE_: case OMP_CLAUSE_INDEPENDENT: case OMP_CLAUSE_WORKER: case OMP_CLAUSE_VECTOR: diff --git gcc/tree-core.h gcc/tree-core.h index abdc2c9..42ad6a0 100644 --- gcc/tree-core.h +++ gcc/tree-core.h @@ -266,6 +266,10 @@ enum omp_clause_code { OpenMP clause: map ({alloc:,to:,from:,tofrom:,}variable-list). */ OMP_CLAUSE_MAP, + /* Internal structure to hold OpenACC cache directive's variable-list. + #pragma acc cache (variable-list). */ + OMP_CLAUSE__CACHE_, + /* OpenACC clause: host (variable_list). */ OMP_CLAUSE_HOST, @@ -292,10 +296,6 @@ enum omp_clause_code { /* OpenACC clause/directive: wait [(integer-expression-list)]. */ OMP_CLAUSE_WAIT, - /* Internal structure to hold OpenACC cache directive's variable-list. - #pragma acc cache (variable-list). */ - OMP_NO_CLAUSE_CACHE, - /* Internal clause: temporary for combined loops expansion. */ OMP_CLAUSE__LOOPTEMP_, diff --git gcc/tree-pretty-print.c gcc/tree-pretty-print.c index f311ed9..d678f36 100644 --- gcc/tree-pretty-print.c +++ gcc/tree-pretty-print.c @@ -347,9 +347,6 @@ dump_omp_clause (pretty_printer *buffer, tree clause, int spc, int flags) case OMP_CLAUSE_USE_DEVICE: name = "use_device"; goto print_remap; - case OMP_NO_CLAUSE_CACHE: - name = "_cache_"; - goto print_remap; print_remap: pp_string (buffer, name); pp_left_paren (buffer); @@ -599,6 +596,12 @@ dump_omp_clause (pretty_printer *buffer, tree clause, int spc, int flags) spc, flags, false); goto print_clause_size; + case OMP_CLAUSE__CACHE_: + pp_string (buffer, "("); + dump_generic_node (buffer, OMP_CLAUSE_DECL (clause), + spc, flags, false); + goto print_clause_size; + case OMP_CLAUSE_NUM_TEAMS: pp_string (buffer, "num_teams("); dump_generic_node (buffer, OMP_CLAUSE_NUM_TEAMS_EXPR (clause), @@ -2548,7 +2551,7 @@ dump_generic_node (pretty_printer *buffer, tree node, int spc, int flags, case OACC_CACHE: pp_string (buffer, "#pragma acc cache"); - dump_omp_clauses (buffer, OACC_CACHE_CLAUSES(node), spc, flags); + dump_omp_clauses (buffer, OACC_CACHE_CLAUSES (node), spc, flags); break; case OMP_PARALLEL: diff --git gcc/tree.c gcc/tree.c index 0475622..f39c63f 100644 --- gcc/tree.c +++ gcc/tree.c @@ -270,6 +270,7 @@ unsigned const char omp_clause_num_ops[] = 2, /* OMP_CLAUSE_FROM */ 2, /* OMP_CLAUSE_TO */ 2, /* OMP_CLAUSE_MAP */ + 2, /* OMP_CLAUSE__CACHE_ */ 1, /* OMP_CLAUSE_HOST */ 1, /* OMP_CLAUSE_OACC_DEVICE */ 1, /* OMP_CLAUSE_DEVICE_RESIDENT */ @@ -277,7 +278,6 @@ unsigned const char omp_clause_num_ops[] = 1, /* OMP_CLAUSE_GANG */ 1, /* OMP_CLAUSE_ASYNC */ 1, /* OMP_CLAUSE_WAIT */ - 1, /* OMP_NO_CLAUSE_CACHE */ 1, /* OMP_CLAUSE__LOOPTEMP_ */ 1, /* OMP_CLAUSE_IF */ 1, /* OMP_CLAUSE_NUM_THREADS */ @@ -329,6 +329,7 @@ const char * const omp_clause_code_name[] = "from", "to", "map", + "_cache_", "host", "device", "device_resident", @@ -336,7 +337,6 @@ const char * const omp_clause_code_name[] = "gang", "async", "wait", - "_cache_", "_looptemp_", "if", "num_threads", @@ -11127,7 +11127,6 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data, case OMP_CLAUSE_GANG: case OMP_CLAUSE_ASYNC: case OMP_CLAUSE_WAIT: - case OMP_NO_CLAUSE_CACHE: case OMP_CLAUSE_WORKER: case OMP_CLAUSE_VECTOR: case OMP_CLAUSE_NUM_GANGS: @@ -11194,6 +11193,7 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data, case OMP_CLAUSE_FROM: case OMP_CLAUSE_TO: case OMP_CLAUSE_MAP: + case OMP_CLAUSE__CACHE_: WALK_SUBTREE (OMP_CLAUSE_DECL (*tp)); WALK_SUBTREE (OMP_CLAUSE_OPERAND (*tp, 1)); WALK_SUBTREE_TAIL (OMP_CLAUSE_CHAIN (*tp)); diff --git gcc/tree.def gcc/tree.def index 871a7fb..f44853a 100644 --- gcc/tree.def +++ gcc/tree.def @@ -1163,8 +1163,9 @@ DEFTREECODE (OACC_ENTER_DATA, "oacc_enter_data", tcc_statement, 1) Operand 0: OACC_EXIT_DATA_CLAUSES: List of clauses. */ DEFTREECODE (OACC_EXIT_DATA, "oacc_exit_data", tcc_statement, 1) -/* OpenACC - #pragma acc cache [clause1 ... clauseN] - Operand 0: OACC_CACHE_CLAUSES: List of clauses. */ +/* OpenACC - #pragma acc cache (variable1 ... variableN) + Operand 0: OACC_CACHE_CLAUSES: List of variables (transformed into + OMP_CLAUSE__CACHE_ clauses). */ DEFTREECODE (OACC_CACHE, "oacc_cache", tcc_statement, 1) /* OpenMP - #pragma omp target update [clause1 ... clauseN] diff --git gcc/tree.h gcc/tree.h index c91e716..e1adbab 100644 --- gcc/tree.h +++ gcc/tree.h @@ -1254,7 +1254,7 @@ extern void protected_set_expr_location (tree, location_t); #define OMP_CLAUSE_SIZE(NODE) \ OMP_CLAUSE_OPERAND (OMP_CLAUSE_RANGE_CHECK (OMP_CLAUSE_CHECK (NODE), \ OMP_CLAUSE_FROM, \ - OMP_CLAUSE_MAP), 1) + OMP_CLAUSE__CACHE_), 1) #define OMP_CLAUSE_CHAIN(NODE) TREE_CHAIN (OMP_CLAUSE_CHECK (NODE)) #define OMP_CLAUSE_DECL(NODE) \