From patchwork Wed Nov 5 21:39:44 2014 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: James Norris X-Patchwork-Id: 407171 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 94BC614003E for ; Thu, 6 Nov 2014 08:40:01 +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 :message-id:date:from:mime-version:to:subject:content-type; q= dns; s=default; b=gbsP0X0vzhekQfzFFefClVIOJJ524dO9PLcpZ4cKoVK9dT 15umz0Z8VFw/Sq9pzr0Ud2tu1YDxuH2hVDGzx+WAt4nIlDLKv68GIUyYSfq+3bHF Src1K6tmsuStFbeKAxyfHMXGeny9ziB2hpgcH3T6bJ9tcRuGrTG2IphzsuZxU= 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 :message-id:date:from:mime-version:to:subject:content-type; s= default; bh=SeojKGBdpG9r361BO9DKNVO74nE=; b=lQ6e1sVSVt4/AD+kTsX9 qWnWVLNdWrxKFCIknwcBJkLBqyBS92+kleCMs9HCa6SM3QY7+D8xcCi6hcQJAjEV 9BdogbV3muO7T5g/x8P3FOz4N2Io8HmNtDiWAxrBI/2UmgvAWNvWwD7RigjEptfN KtIp/23Zik/96o/vDjyKfBw= Received: (qmail 20129 invoked by alias); 5 Nov 2014 21:39:53 -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 20120 invoked by uid 89); 5 Nov 2014 21:39:52 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=0.5 required=5.0 tests=AWL, BAYES_00, RCVD_IN_DNSWL_NONE, UNWANTED_LANGUAGE_BODY 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 21:39:48 +0000 Received: from svr-orw-fem-03.mgc.mentorg.com ([147.34.97.39]) by relay1.mentorg.com with esmtp id 1Xm8Ir-0006oo-Bj from James_Norris@mentor.com for gcc-patches@gcc.gnu.org; Wed, 05 Nov 2014 13:39:45 -0800 Received: from [172.30.36.100] (147.34.91.1) by svr-orw-fem-03.mgc.mentorg.com (147.34.97.39) with Microsoft SMTP Server id 14.3.181.6; Wed, 5 Nov 2014 13:39:44 -0800 Message-ID: <545A9920.8060709@codesourcery.com> Date: Wed, 5 Nov 2014 15:39:44 -0600 From: James Norris User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:31.0) Gecko/20100101 Thunderbird/31.2.0 MIME-Version: 1.0 To: Subject: [PATCH] OpenACC for C front end Hi! This patch represents the changes for OpenACC 2.0 in the C front-end. At present these files will not compile as the changes for the middle end are not present. OK to commit? Thanks, Jim => c/ChangeLog 2014-11-05 James Norris Cesar Philippidis Thomas Schwinge Ilmir Usmanov * c-typeck.c (c_finish_oacc_parallel, c_finish_oacc_kernels, c_finish_oacc_data): New functions. (handle_omp_array_sections, c_finish_omp_clauses): Handle OMP_CLAUSE_MAP_FORCE_DEVICEPTR, OMP_CLAUSE_NUM_GANGS, OMP_CLAUSE_NUM_WORKERS, OMP_CLAUSE_VECTOR_LENGTH, OMP_CLAUSE_ASYNC, and OMP_CLAUSE_WAIT. (c_finish_omp_clauses): Handle PRAGMA_OMP_CLAUSE_NUM_GANGS, PRAGMA_OMP_CLAUSE_NUM_WORKERS, and PRAGMA_OMP_CLAUSE_VECTOR_LENGTH. * c-tree.h (c_finish_oacc_parallel, c_finish_oacc_kernels, c_finish_oacc_data): New prototypes. * c-parser.c (c_parser_pragma): Handle PRAGMA_OACC_ENTER_DATA, PRAGMA_OACC_EXIT_DATA, and PRAGMA_OACC_UPDATE. (c_parser_omp_clause_name): Handle OpenACC clauses. (c_parser_oacc_wait_list, c_parser_oacc_data_clause, c_parser_oacc_data_clause_deviceptr, c_parser_omp_clause_num_gangs, c_parser_omp_clause_num_workers, c_parser_oacc_clause_async, c_parser_oacc_clause_wait, c_parser_omp_clause_vector_length, c_parser_oacc_all_clauses, c_parser_oacc_data, c_parser_oacc_kernels, c_parser_oacc_enter_exit_data, c_parser_oacc_loop, c_parser_oacc_parallel, c_parser_oacc_update, c_parser_oacc_wait): New functions. (c_parser_omp_construct): Handle PRAGMA_OMP_DATA, PRAGMA_OACC_KERNELS, PRAGMA_OACC_LOOP, PRAGMA_OACC_PARALLEL and PRAMGA_OACC_WAIT. diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index d316216..df4b4cb 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -1239,10 +1239,15 @@ static vec *c_parser_expr_list (c_parser *, bool, bool, vec **, location_t *, tree *, vec *, unsigned int * = NULL); +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_oacc_enter_exit_data (c_parser *, bool); +static void c_parser_oacc_update (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 *); @@ -4482,6 +4487,14 @@ c_parser_initval (c_parser *parser, struct c_expr *after, Although they are erroneous if the labels declared aren't defined, is it useful for the syntax to be this way? + OpenACC: + + block-item: + openacc-directive + + openacc-directive: + update-directive + OpenMP: block-item: @@ -4828,6 +4841,29 @@ c_parser_label (c_parser *parser) @throw expression ; @throw ; + OpenACC: + + statement: + openacc-construct + + openacc-construct: + parallel-construct + kernels-construct + data-construct + loop-construct + + parallel-construct: + parallel-directive structured-block + + kernels-construct: + kernels-directive structured-block + + data-construct: + data-directive structured-block + + loop-construct: + loop-directive structured-block + OpenMP: statement: @@ -9509,6 +9545,25 @@ c_parser_pragma (c_parser *parser, enum pragma_context context) switch (id) { + case PRAGMA_OACC_ENTER_DATA: + c_parser_oacc_enter_exit_data (parser, true); + return false; + + case PRAGMA_OACC_EXIT_DATA: + c_parser_oacc_enter_exit_data (parser, false); + return false; + + case PRAGMA_OACC_UPDATE: + if (context != pragma_compound) + { + if (context == pragma_stmt) + c_parser_error (parser, "%<#pragma acc update%> may only be " + "used in compound statements"); + goto bad_stmt; + } + c_parser_oacc_update (parser); + return false; + case PRAGMA_OMP_BARRIER: if (context != pragma_compound) { @@ -9711,7 +9766,7 @@ c_parser_pragma_pch_preprocess (c_parser *parser) c_common_pch_pragma (parse_in, TREE_STRING_POINTER (name)); } -/* OpenMP 2.5 / 3.0 / 3.1 / 4.0 parsing routines. */ +/* OpenACC and OpenMP parsing routines. */ /* Returns name of the next clause. If the clause is not recognized PRAGMA_OMP_CLAUSE_NONE is returned and @@ -9738,20 +9793,32 @@ c_parser_omp_clause_name (c_parser *parser) case 'a': if (!strcmp ("aligned", p)) result = PRAGMA_OMP_CLAUSE_ALIGNED; + else if (!strcmp ("async", p)) + result = PRAGMA_OMP_CLAUSE_ASYNC; break; case 'c': if (!strcmp ("collapse", p)) result = PRAGMA_OMP_CLAUSE_COLLAPSE; + else if (!strcmp ("copy", p)) + result = PRAGMA_OMP_CLAUSE_COPY; else if (!strcmp ("copyin", p)) result = PRAGMA_OMP_CLAUSE_COPYIN; + else if (!strcmp ("copyout", p)) + result = PRAGMA_OMP_CLAUSE_COPYOUT; else if (!strcmp ("copyprivate", p)) result = PRAGMA_OMP_CLAUSE_COPYPRIVATE; + else if (!strcmp ("create", p)) + result = PRAGMA_OMP_CLAUSE_CREATE; break; case 'd': - if (!strcmp ("depend", p)) + if (!strcmp ("delete", p)) + result = PRAGMA_OMP_CLAUSE_DELETE; + else if (!strcmp ("depend", p)) result = PRAGMA_OMP_CLAUSE_DEPEND; else if (!strcmp ("device", p)) result = PRAGMA_OMP_CLAUSE_DEVICE; + else if (!strcmp ("deviceptr", p)) + result = PRAGMA_OMP_CLAUSE_DEVICEPTR; else if (!strcmp ("dist_schedule", p)) result = PRAGMA_OMP_CLAUSE_DIST_SCHEDULE; break; @@ -9763,6 +9830,10 @@ c_parser_omp_clause_name (c_parser *parser) else if (!strcmp ("from", p)) result = PRAGMA_OMP_CLAUSE_FROM; break; + case 'h': + if (!strcmp ("host", p)) + result = PRAGMA_OMP_CLAUSE_SELF; + break; case 'i': if (!strcmp ("inbranch", p)) result = PRAGMA_OMP_CLAUSE_INBRANCH; @@ -9786,10 +9857,14 @@ c_parser_omp_clause_name (c_parser *parser) result = PRAGMA_OMP_CLAUSE_NOTINBRANCH; else if (!strcmp ("nowait", p)) result = PRAGMA_OMP_CLAUSE_NOWAIT; + else if (!strcmp ("num_gangs", p)) + result = PRAGMA_OMP_CLAUSE_NUM_GANGS; else if (!strcmp ("num_teams", p)) result = PRAGMA_OMP_CLAUSE_NUM_TEAMS; else if (!strcmp ("num_threads", p)) result = PRAGMA_OMP_CLAUSE_NUM_THREADS; + else if (!strcmp ("num_workers", p)) + result = PRAGMA_OMP_CLAUSE_NUM_WORKERS; else if (flag_cilkplus && !strcmp ("nomask", p)) result = PRAGMA_CILK_CLAUSE_NOMASK; break; @@ -9800,6 +9875,20 @@ c_parser_omp_clause_name (c_parser *parser) case 'p': if (!strcmp ("parallel", p)) result = PRAGMA_OMP_CLAUSE_PARALLEL; + else if (!strcmp ("present", p)) + result = PRAGMA_OMP_CLAUSE_PRESENT; + else if (!strcmp ("present_or_copy", p) + || !strcmp ("pcopy", p)) + result = PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY; + else if (!strcmp ("present_or_copyin", p) + || !strcmp ("pcopyin", p)) + result = PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN; + else if (!strcmp ("present_or_copyout", p) + || !strcmp ("pcopyout", p)) + result = PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT; + else if (!strcmp ("present_or_create", p) + || !strcmp ("pcreate", p)) + result = PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE; else if (!strcmp ("private", p)) result = PRAGMA_OMP_CLAUSE_PRIVATE; else if (!strcmp ("proc_bind", p)) @@ -9820,6 +9909,8 @@ c_parser_omp_clause_name (c_parser *parser) result = PRAGMA_OMP_CLAUSE_SHARED; else if (!strcmp ("simdlen", p)) result = PRAGMA_OMP_CLAUSE_SIMDLEN; + else if (!strcmp ("self", p)) + result = PRAGMA_OMP_CLAUSE_SELF; break; case 't': if (!strcmp ("taskgroup", p)) @@ -9836,9 +9927,15 @@ c_parser_omp_clause_name (c_parser *parser) result = PRAGMA_OMP_CLAUSE_UNTIED; break; case 'v': - if (flag_cilkplus && !strcmp ("vectorlength", p)) + if (!strcmp ("vector_length", p)) + result = PRAGMA_OMP_CLAUSE_VECTOR_LENGTH; + else if (flag_cilkplus && !strcmp ("vectorlength", p)) result = PRAGMA_CILK_CLAUSE_VECTORLENGTH; break; + case 'w': + if (!strcmp ("wait", p)) + result = PRAGMA_OMP_CLAUSE_WAIT; + break; } } @@ -9865,7 +9962,57 @@ check_no_duplicate_clause (tree clauses, enum omp_clause_code code, } } -/* OpenMP 2.5: +/* OpenACC 2.0 + Parse wait clause or wait directive parameters. */ + +static tree +c_parser_oacc_wait_list (c_parser *parser, location_t clause_loc, tree list) +{ + vec *args; + tree t, args_tree; + + if (!c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>")) + return list; + + args = c_parser_expr_list (parser, false, true, NULL, NULL, NULL, NULL); + + if (args->length () == 0) + { + c_parser_error (parser, "expected integer expression before ')'"); + release_tree_vector (args); + return list; + } + + args_tree = build_tree_list_vec (args); + + for (t = args_tree; t; t = TREE_CHAIN (t)) + { + tree targ = TREE_VALUE (t); + + if (targ != error_mark_node) + { + if (!INTEGRAL_TYPE_P (TREE_TYPE (targ))) + { + c_parser_error (parser, "expression must be integral"); + targ = error_mark_node; + } + else + { + tree c = build_omp_clause (clause_loc, OMP_CLAUSE_WAIT); + + OMP_CLAUSE_DECL (c) = targ; + OMP_CLAUSE_CHAIN (c) = list; + list = c; + } + } + } + + release_tree_vector (args); + c_parser_require (parser, CPP_CLOSE_PAREN, "expected %<)%>"); + return list; +} + +/* OpenACC 2.0, OpenMP 2.5: variable-list: identifier variable-list , identifier @@ -9972,7 +10119,7 @@ c_parser_omp_variable_list (c_parser *parser, } /* Similarly, but expect leading and trailing parenthesis. This is a very - common case for omp clauses. */ + common case for OpenACC and OpenMP clauses. */ static tree c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind, @@ -9989,7 +10136,121 @@ c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind, return list; } -/* OpenMP 3.0: +/* OpenACC 2.0: + copy ( variable-list ) + copyin ( variable-list ) + copyout ( variable-list ) + create ( variable-list ) + delete ( variable-list ) + present ( variable-list ) + present_or_copy ( variable-list ) + pcopy ( variable-list ) + present_or_copyin ( variable-list ) + pcopyin ( variable-list ) + present_or_copyout ( variable-list ) + pcopyout ( variable-list ) + present_or_create ( variable-list ) + pcreate ( variable-list ) */ + +static tree +c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind, + tree list) +{ + enum omp_clause_map_kind kind; + switch (c_kind) + { + default: + gcc_unreachable (); + case PRAGMA_OMP_CLAUSE_COPY: + kind = OMP_CLAUSE_MAP_FORCE_TOFROM; + break; + case PRAGMA_OMP_CLAUSE_COPYIN: + kind = OMP_CLAUSE_MAP_FORCE_TO; + break; + case PRAGMA_OMP_CLAUSE_COPYOUT: + kind = OMP_CLAUSE_MAP_FORCE_FROM; + break; + case PRAGMA_OMP_CLAUSE_CREATE: + kind = OMP_CLAUSE_MAP_FORCE_ALLOC; + break; + case PRAGMA_OMP_CLAUSE_DELETE: + kind = OMP_CLAUSE_MAP_FORCE_DEALLOC; + break; + case PRAGMA_OMP_CLAUSE_DEVICE: + kind = OMP_CLAUSE_MAP_FORCE_TO; + break; + case PRAGMA_OMP_CLAUSE_HOST: + kind = OMP_CLAUSE_MAP_FORCE_FROM; + break; + case PRAGMA_OMP_CLAUSE_PRESENT: + kind = OMP_CLAUSE_MAP_FORCE_PRESENT; + break; + case PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY: + kind = OMP_CLAUSE_MAP_TOFROM; + break; + case PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN: + kind = OMP_CLAUSE_MAP_TO; + break; + case PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT: + kind = OMP_CLAUSE_MAP_FROM; + break; + case PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE: + kind = OMP_CLAUSE_MAP_ALLOC; + break; + case PRAGMA_OMP_CLAUSE_SELF: + kind = OMP_CLAUSE_MAP_FORCE_FROM; + break; + } + tree nl, c; + nl = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list); + + for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c)) + OMP_CLAUSE_MAP_KIND (c) = kind; + + return nl; +} + +/* OpenACC 2.0: + deviceptr ( variable-list ) */ + +static tree +c_parser_oacc_data_clause_deviceptr (c_parser *parser, tree list) +{ + location_t loc = c_parser_peek_token (parser)->location; + tree vars, t; + + /* Can't use OMP_CLAUSE_MAP here (that is, can't use the generic + c_parser_oacc_data_clause), as for PRAGMA_OMP_CLAUSE_DEVICEPTR, + variable-list must only allow for pointer variables. */ + vars = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_ERROR, NULL); + for (t = vars; t && t; t = TREE_CHAIN (t)) + { + tree v = TREE_PURPOSE (t); + + /* FIXME diagnostics: Ideally we should keep individual + locations for all the variables in the var list to make the + following errors more precise. Perhaps + c_parser_omp_var_list_parens() should construct a list of + locations to go along with the var list. */ + + if (TREE_CODE (v) != VAR_DECL) + error_at (loc, "%qD is not a variable", v); + else if (TREE_TYPE (v) == error_mark_node) + ; + else if (!POINTER_TYPE_P (TREE_TYPE (v))) + error_at (loc, "%qD is not a pointer variable", v); + + tree u = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_MAP_KIND (u) = OMP_CLAUSE_MAP_FORCE_DEVICEPTR; + OMP_CLAUSE_DECL (u) = v; + OMP_CLAUSE_CHAIN (u) = list; + list = u; + } + + return list; +} + +/* OpenACC 2.0, OpenMP 3.0: collapse ( constant-expression ) */ static tree @@ -10132,7 +10393,7 @@ c_parser_omp_clause_final (c_parser *parser, tree list) return list; } -/* OpenMP 2.5: +/* OpenACC, OpenMP 2.5: if ( expression ) */ static tree @@ -10200,6 +10461,51 @@ c_parser_omp_clause_nowait (c_parser *parser ATTRIBUTE_UNUSED, tree list) return c; } +/* OpenACC: + num_gangs ( expression ) */ + +static tree +c_parser_omp_clause_num_gangs (c_parser *parser, tree list) +{ + location_t num_gangs_loc = c_parser_peek_token (parser)->location; + if (c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>")) + { + location_t expr_loc = c_parser_peek_token (parser)->location; + tree c, t = c_parser_expression (parser).value; + mark_exp_read (t); + t = c_fully_fold (t, false, NULL); + + c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, "expected %<)%>"); + + if (!INTEGRAL_TYPE_P (TREE_TYPE (t))) + { + c_parser_error (parser, "expected integer expression"); + return list; + } + + /* Attempt to statically determine when the number isn't positive. */ + c = fold_build2_loc (expr_loc, LE_EXPR, boolean_type_node, t, + build_int_cst (TREE_TYPE (t), 0)); + if (CAN_HAVE_LOCATION_P (c)) + SET_EXPR_LOCATION (c, expr_loc); + if (c == boolean_true_node) + { + warning_at (expr_loc, 0, + "% value must be positive"); + t = integer_one_node; + } + + check_no_duplicate_clause (list, OMP_CLAUSE_NUM_GANGS, "num_gangs"); + + c = build_omp_clause (num_gangs_loc, OMP_CLAUSE_NUM_GANGS); + OMP_CLAUSE_NUM_GANGS_EXPR (c) = t; + OMP_CLAUSE_CHAIN (c) = list; + list = c; + } + + return list; +} + /* OpenMP 2.5: num_threads ( expression ) */ @@ -10245,6 +10551,103 @@ c_parser_omp_clause_num_threads (c_parser *parser, tree list) return list; } +/* OpenACC: + num_workers ( expression ) */ + +static tree +c_parser_omp_clause_num_workers (c_parser *parser, tree list) +{ + location_t num_workers_loc = c_parser_peek_token (parser)->location; + if (c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>")) + { + location_t expr_loc = c_parser_peek_token (parser)->location; + tree c, t = c_parser_expression (parser).value; + mark_exp_read (t); + t = c_fully_fold (t, false, NULL); + + c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, "expected %<)%>"); + + if (!INTEGRAL_TYPE_P (TREE_TYPE (t))) + { + c_parser_error (parser, "expected integer expression"); + return list; + } + + /* Attempt to statically determine when the number isn't positive. */ + c = fold_build2_loc (expr_loc, LE_EXPR, boolean_type_node, t, + build_int_cst (TREE_TYPE (t), 0)); + if (CAN_HAVE_LOCATION_P (c)) + SET_EXPR_LOCATION (c, expr_loc); + if (c == boolean_true_node) + { + warning_at (expr_loc, 0, + "% value must be positive"); + t = integer_one_node; + } + + check_no_duplicate_clause (list, OMP_CLAUSE_NUM_WORKERS, "num_workers"); + + c = build_omp_clause (num_workers_loc, OMP_CLAUSE_NUM_WORKERS); + OMP_CLAUSE_NUM_WORKERS_EXPR (c) = t; + OMP_CLAUSE_CHAIN (c) = list; + list = c; + } + + return list; +} + +/* OpenACC: + async [( int-expr )] */ + +static tree +c_parser_oacc_clause_async (c_parser *parser, tree list) +{ + tree c, t; + location_t loc = c_parser_peek_token (parser)->location; + + /* TODO XXX: FIX -1 (acc_async_noval). */ + t = build_int_cst (integer_type_node, -1); + + if (c_parser_peek_token (parser)->type == CPP_OPEN_PAREN) + { + c_parser_consume_token (parser); + + t = c_parser_expression (parser).value; + if (!INTEGRAL_TYPE_P (TREE_TYPE (t))) + c_parser_error (parser, "expected integer expression"); + else if (t == error_mark_node + || !c_parser_require (parser, CPP_CLOSE_PAREN, "expected %<)%>")) + return list; + } + else + { + t = c_fully_fold (t, false, NULL); + } + + check_no_duplicate_clause (list, OMP_CLAUSE_ASYNC, "async"); + + c = build_omp_clause (loc, OMP_CLAUSE_ASYNC); + OMP_CLAUSE_ASYNC_EXPR (c) = t; + OMP_CLAUSE_CHAIN (c) = list; + list = c; + + return list; +} + +/* OpenACC: + wait ( int-expr-list ) */ + +static tree +c_parser_oacc_clause_wait (c_parser *parser, tree list) +{ + location_t clause_loc = c_parser_peek_token (parser)->location; + + if (c_parser_peek_token (parser)->type == CPP_OPEN_PAREN) + list = c_parser_oacc_wait_list (parser, clause_loc, list); + + return list; +} + /* OpenMP 2.5: ordered */ @@ -10496,33 +10899,78 @@ c_parser_omp_clause_untied (c_parser *parser ATTRIBUTE_UNUSED, tree list) return c; } -/* OpenMP 4.0: - inbranch - notinbranch */ +/* OpenACC: + vector_length ( expression ) */ static tree -c_parser_omp_clause_branch (c_parser *parser ATTRIBUTE_UNUSED, - enum omp_clause_code code, tree list) +c_parser_omp_clause_vector_length (c_parser *parser, tree list) { - check_no_duplicate_clause (list, code, omp_clause_code_name[code]); - - tree c = build_omp_clause (c_parser_peek_token (parser)->location, code); - OMP_CLAUSE_CHAIN (c) = list; + location_t vector_length_loc = c_parser_peek_token (parser)->location; + if (c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>")) + { + location_t expr_loc = c_parser_peek_token (parser)->location; + tree c, t = c_parser_expression (parser).value; + mark_exp_read (t); + t = c_fully_fold (t, false, NULL); - return c; -} + c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, "expected %<)%>"); -/* OpenMP 4.0: - parallel - for - sections - taskgroup */ + if (!INTEGRAL_TYPE_P (TREE_TYPE (t))) + { + c_parser_error (parser, "expected integer expression"); + return list; + } -static tree -c_parser_omp_clause_cancelkind (c_parser *parser ATTRIBUTE_UNUSED, - enum omp_clause_code code, tree list) -{ - tree c = build_omp_clause (c_parser_peek_token (parser)->location, code); + /* Attempt to statically determine when the number isn't positive. */ + c = fold_build2_loc (expr_loc, LE_EXPR, boolean_type_node, t, + build_int_cst (TREE_TYPE (t), 0)); + if (CAN_HAVE_LOCATION_P (c)) + SET_EXPR_LOCATION (c, expr_loc); + if (c == boolean_true_node) + { + warning_at (expr_loc, 0, + "% value must be positive"); + t = integer_one_node; + } + + check_no_duplicate_clause (list, OMP_CLAUSE_VECTOR_LENGTH, "vector_length"); + + c = build_omp_clause (vector_length_loc, OMP_CLAUSE_VECTOR_LENGTH); + OMP_CLAUSE_VECTOR_LENGTH_EXPR (c) = t; + OMP_CLAUSE_CHAIN (c) = list; + list = c; + } + + return list; +} + +/* OpenMP 4.0: + inbranch + notinbranch */ + +static tree +c_parser_omp_clause_branch (c_parser *parser ATTRIBUTE_UNUSED, + enum omp_clause_code code, tree list) +{ + check_no_duplicate_clause (list, code, omp_clause_code_name[code]); + + tree c = build_omp_clause (c_parser_peek_token (parser)->location, code); + OMP_CLAUSE_CHAIN (c) = list; + + return c; +} + +/* OpenMP 4.0: + parallel + for + sections + taskgroup */ + +static tree +c_parser_omp_clause_cancelkind (c_parser *parser ATTRIBUTE_UNUSED, + enum omp_clause_code code, tree list) +{ + tree c = build_omp_clause (c_parser_peek_token (parser)->location, code); OMP_CLAUSE_CHAIN (c) = list; return c; @@ -11032,6 +11480,144 @@ c_parser_omp_clause_uniform (c_parser *parser, tree list) return list; } +/* Parse all OpenACC clauses. The set clauses allowed by the directive + is a bitmask in MASK. Return the list of clauses found. */ + +static tree +c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, + const char *where, bool finish_p = true) +{ + tree clauses = NULL; + bool first = true; + + while (c_parser_next_token_is_not (parser, CPP_PRAGMA_EOL)) + { + location_t here; + pragma_omp_clause c_kind; + const char *c_name; + tree prev = clauses; + + if (!first && c_parser_next_token_is (parser, CPP_COMMA)) + c_parser_consume_token (parser); + + here = c_parser_peek_token (parser)->location; + c_kind = c_parser_omp_clause_name (parser); + + switch (c_kind) + { + case PRAGMA_OMP_CLAUSE_ASYNC: + clauses = c_parser_oacc_clause_async (parser, clauses); + c_name = "async"; + break; + case PRAGMA_OMP_CLAUSE_COLLAPSE: + clauses = c_parser_omp_clause_collapse (parser, clauses); + c_name = "collapse"; + break; + case PRAGMA_OMP_CLAUSE_COPY: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "copy"; + break; + case PRAGMA_OMP_CLAUSE_COPYIN: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "copyin"; + break; + case PRAGMA_OMP_CLAUSE_COPYOUT: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "copyout"; + break; + case PRAGMA_OMP_CLAUSE_CREATE: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "create"; + break; + case PRAGMA_OMP_CLAUSE_DELETE: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "delete"; + break; + case PRAGMA_OMP_CLAUSE_DEVICE: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "device"; + break; + case PRAGMA_OMP_CLAUSE_DEVICEPTR: + clauses = c_parser_oacc_data_clause_deviceptr (parser, clauses); + c_name = "deviceptr"; + break; + case PRAGMA_OMP_CLAUSE_HOST: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "host"; + break; + case PRAGMA_OMP_CLAUSE_IF: + clauses = c_parser_omp_clause_if (parser, clauses); + c_name = "if"; + break; + case PRAGMA_OMP_CLAUSE_NUM_GANGS: + clauses = c_parser_omp_clause_num_gangs (parser, clauses); + c_name = "num_gangs"; + break; + case PRAGMA_OMP_CLAUSE_NUM_WORKERS: + clauses = c_parser_omp_clause_num_workers (parser, clauses); + c_name = "num_workers"; + break; + case PRAGMA_OMP_CLAUSE_PRESENT: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "present"; + break; + case PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "present_or_copy"; + break; + case PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "present_or_copyin"; + break; + case PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "present_or_copyout"; + break; + case PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "present_or_create"; + break; + case PRAGMA_OMP_CLAUSE_REDUCTION: + clauses = c_parser_omp_clause_reduction (parser, clauses); + c_name = "reduction"; + break; + case PRAGMA_OMP_CLAUSE_SELF: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "self"; + break; + case PRAGMA_OMP_CLAUSE_VECTOR_LENGTH: + clauses = c_parser_omp_clause_vector_length (parser, clauses); + c_name = "vector_length"; + break; + case PRAGMA_OMP_CLAUSE_WAIT: + clauses = c_parser_oacc_clause_wait (parser, clauses); + c_name = "wait"; + break; + default: + c_parser_error (parser, "expected clause"); + goto saw_error; + } + + first = false; + + if (((mask >> c_kind) & 1) == 0 && !parser->error) + { + /* Remove the invalid clause(s) from the list to avoid + confusing the rest of the compiler. */ + clauses = prev; + error_at (here, "%qs is not valid for %qs", c_name, where); + } + } + + saw_error: + c_parser_skip_to_pragma_eol (parser); + + if (finish_p) + return c_finish_omp_clauses (clauses); + + return clauses; +} + /* Parse all OpenMP clauses. The set clauses allowed by the directive is a bitmask in MASK. Return the list of clauses found; the result of clause default goes in *pdefault. */ @@ -11262,7 +11848,7 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask, return clauses; } -/* OpenMP 2.5: +/* OpenACC 2.0, OpenMP 2.5: structured-block: statement @@ -11278,6 +11864,326 @@ c_parser_omp_structured_block (c_parser *parser) return pop_stmt_list (stmt); } +/* OpenACC 2.0: + # pragma acc data oacc-data-clause[optseq] new-line + structured-block + + LOC is the location of the #pragma token. +*/ + +#define OACC_DATA_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPY) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYIN) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYOUT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_CREATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICEPTR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE) ) + +static tree +c_parser_oacc_data (location_t loc, c_parser *parser) +{ + tree stmt, clauses, block; + + clauses = c_parser_oacc_all_clauses (parser, OACC_DATA_CLAUSE_MASK, + "#pragma acc data"); + + block = c_begin_omp_parallel (); + add_stmt (c_parser_omp_structured_block (parser)); + + stmt = c_finish_oacc_data (loc, clauses, block); + + return stmt; +} + +/* OpenACC 2.0: + # pragma acc kernels oacc-kernels-clause[optseq] new-line + structured-block + + LOC is the location of the #pragma token. +*/ + +#define OACC_KERNELS_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_ASYNC) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPY) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYIN) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYOUT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_CREATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICEPTR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_WAIT) ) + +static tree +c_parser_oacc_kernels (location_t loc, c_parser *parser, char *p_name) +{ + tree stmt, clauses = NULL_TREE, block; + + strcat (p_name, " kernels"); + + if (c_parser_next_token_is (parser, CPP_NAME)) + { + const char *p = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value); + if (strcmp (p, "loop") == 0) + { + c_parser_consume_token (parser); + block = c_begin_omp_parallel (); + c_parser_oacc_loop (loc, parser, p_name); + stmt = c_finish_oacc_kernels (loc, clauses, block); + OACC_KERNELS_COMBINED (stmt) = 1; + return stmt; + } + } + + clauses = c_parser_oacc_all_clauses (parser, OACC_KERNELS_CLAUSE_MASK, + p_name); + + block = c_begin_omp_parallel (); + add_stmt (c_parser_omp_structured_block (parser)); + + stmt = c_finish_oacc_kernels (loc, clauses, block); + + return stmt; +} + +/* OpenACC 2.0: + # pragma acc enter data oacc-enter-data-clause[optseq] new-line + + or + + # pragma acc exit data oacc-exit-data-clause[optseq] new-line + + + LOC is the location of the #pragma token. +*/ + +#define OACC_ENTER_DATA_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_ASYNC) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYIN) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_CREATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_WAIT) ) + +#define OACC_EXIT_DATA_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_ASYNC) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYOUT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DELETE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_WAIT) ) + +static void +c_parser_oacc_enter_exit_data (c_parser *parser, bool enter) +{ + location_t loc = c_parser_peek_token (parser)->location; + tree clauses, stmt; + + c_parser_consume_pragma (parser); + + if (!c_parser_next_token_is (parser, CPP_NAME)) + { + c_parser_error (parser, enter + ? "expected % in %<#pragma acc enter data%>" + : "expected % in %<#pragma acc exit data%>"); + c_parser_skip_to_pragma_eol (parser); + return; + } + + const char *p = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value); + if (strcmp (p, "data") != 0) + { + c_parser_error (parser, "invalid pragma"); + c_parser_skip_to_pragma_eol (parser); + return; + } + + c_parser_consume_token (parser); + + if (enter) + clauses = c_parser_oacc_all_clauses (parser, OACC_ENTER_DATA_CLAUSE_MASK, + "#pragma acc enter data"); + else + clauses = c_parser_oacc_all_clauses (parser, OACC_EXIT_DATA_CLAUSE_MASK, + "#pragma acc exit data"); + + if (find_omp_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE) + { + error_at (loc, enter + ? "%<#pragma acc enter data%> has no data movement clause" + : "%<#pragma acc exit data%> has no data movement clause"); + return; + } + + stmt = enter ? make_node (OACC_ENTER_DATA) : make_node (OACC_EXIT_DATA);; + TREE_TYPE (stmt) = void_type_node; + if (enter) + OACC_ENTER_DATA_CLAUSES (stmt) = clauses; + else + OACC_EXIT_DATA_CLAUSES (stmt) = clauses; + SET_EXPR_LOCATION (stmt, loc); + add_stmt (stmt); +} + + +/* 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_COLLAPSE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_REDUCTION) ) + +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 + + LOC is the location of the #pragma token. +*/ + +#define OACC_PARALLEL_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_ASYNC) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPY) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYIN) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYOUT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_CREATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICEPTR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NUM_GANGS) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NUM_WORKERS) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_REDUCTION) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_VECTOR_LENGTH) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_WAIT) ) + +static tree +c_parser_oacc_parallel (location_t loc, c_parser *parser, char *p_name) +{ + tree stmt, clauses = NULL_TREE, block; + + strcat (p_name, " parallel"); + + if (c_parser_next_token_is (parser, CPP_NAME)) + { + const char *p = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value); + if (strcmp (p, "loop") == 0) + { + c_parser_consume_token (parser); + block = c_begin_omp_parallel (); + c_parser_oacc_loop (loc, parser, p_name); + stmt = c_finish_oacc_parallel (loc, clauses, block); + OACC_PARALLEL_COMBINED (stmt) = 1; + return stmt; + } + } + + clauses = c_parser_oacc_all_clauses (parser, OACC_PARALLEL_CLAUSE_MASK, + p_name); + + block = c_begin_omp_parallel (); + add_stmt (c_parser_omp_structured_block (parser)); + + stmt = c_finish_oacc_parallel (loc, clauses, block); + + return stmt; +} + +/* OpenACC 2.0: + # pragma acc update oacc-update-clause[optseq] new-line +*/ + +#define OACC_UPDATE_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_ASYNC) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_HOST) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_SELF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_WAIT) ) + +static void +c_parser_oacc_update (c_parser *parser) +{ + location_t loc = c_parser_peek_token (parser)->location; + + c_parser_consume_pragma (parser); + + tree clauses = c_parser_oacc_all_clauses (parser, OACC_UPDATE_CLAUSE_MASK, + "#pragma acc update"); + if (find_omp_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE) + { + error_at (loc, + "%<#pragma acc update%> must contain at least one " + "% or % clause"); + return; + } + + if (parser->error) + return; + + tree stmt = make_node (OACC_UPDATE); + TREE_TYPE (stmt) = void_type_node; + OACC_UPDATE_CLAUSES (stmt) = clauses; + SET_EXPR_LOCATION (stmt, loc); + add_stmt (stmt); +} + +/* OpenACC 2.0: + # pragma acc wait [(intseq)] oacc-wait-clause[optseq] new-line + + LOC is the location of the #pragma token. +*/ + +#define OACC_WAIT_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_ASYNC) ) + +static tree +c_parser_oacc_wait (location_t loc, c_parser *parser, char *p_name) +{ + tree clauses, list = NULL_TREE, stmt = NULL_TREE; + + if (c_parser_peek_token (parser)->type == CPP_OPEN_PAREN) + list = c_parser_oacc_wait_list (parser, loc, list); + + strcpy (p_name, " wait"); + clauses = c_parser_oacc_all_clauses (parser, OACC_WAIT_CLAUSE_MASK, p_name); + stmt = c_finish_oacc_wait (loc, list, clauses); + + return stmt; +} + /* OpenMP 2.5: # pragma omp atomic new-line expression-stmt @@ -11754,10 +12660,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, @@ -12010,6 +12917,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 @@ -13599,6 +14507,25 @@ c_parser_omp_construct (c_parser *parser) switch (p_kind) { + case PRAGMA_OACC_DATA: + stmt = c_parser_oacc_data (loc, parser); + break; + case PRAGMA_OACC_KERNELS: + strcpy (p_name, "#pragma acc"); + stmt = c_parser_oacc_kernels (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_OACC_WAIT: + strcpy (p_name, "#pragma wait"); + stmt = c_parser_oacc_wait (loc, parser, p_name); + break; case PRAGMA_OMP_ATOMIC: c_parser_omp_atomic (loc, parser); return; diff --git a/gcc/c/c-tree.h b/gcc/c/c-tree.h index f7e723b..bcfec28 100644 --- a/gcc/c/c-tree.h +++ b/gcc/c/c-tree.h @@ -640,6 +640,9 @@ extern tree c_finish_bc_stmt (location_t, tree *, bool); extern tree c_finish_goto_label (location_t, tree); extern tree c_finish_goto_ptr (location_t, tree); extern tree c_expr_to_decl (tree, bool *, bool *); +extern tree c_finish_oacc_parallel (location_t, tree, tree); +extern tree c_finish_oacc_kernels (location_t, tree, tree); +extern tree c_finish_oacc_data (location_t, tree, tree); extern tree c_begin_omp_parallel (void); extern tree c_finish_omp_parallel (location_t, tree, tree); extern tree c_begin_omp_task (void); diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index 79dbc3d..76503e4 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -11230,6 +11230,63 @@ c_expr_to_decl (tree expr, bool *tc ATTRIBUTE_UNUSED, bool *se) return expr; } +/* Generate OACC_PARALLEL, with CLAUSES and BLOCK as its compound + statement. LOC is the location of the OACC_PARALLEL. */ + +tree +c_finish_oacc_parallel (location_t loc, tree clauses, tree block) +{ + tree stmt; + + block = c_end_compound_stmt (loc, block, true); + + stmt = make_node (OACC_PARALLEL); + TREE_TYPE (stmt) = void_type_node; + OACC_PARALLEL_CLAUSES (stmt) = clauses; + OACC_PARALLEL_BODY (stmt) = block; + SET_EXPR_LOCATION (stmt, loc); + + return add_stmt (stmt); +} + +/* Generate OACC_KERNELS, with CLAUSES and BLOCK as its compound + statement. LOC is the location of the OACC_KERNELS. */ + +tree +c_finish_oacc_kernels (location_t loc, tree clauses, tree block) +{ + tree stmt; + + block = c_end_compound_stmt (loc, block, true); + + stmt = make_node (OACC_KERNELS); + TREE_TYPE (stmt) = void_type_node; + OACC_KERNELS_CLAUSES (stmt) = clauses; + OACC_KERNELS_BODY (stmt) = block; + SET_EXPR_LOCATION (stmt, loc); + + return add_stmt (stmt); +} + +/* Generate OACC_DATA, with CLAUSES and BLOCK as its compound + statement. LOC is the location of the OACC_DATA. */ + +tree +c_finish_oacc_data (location_t loc, tree clauses, tree block) +{ + tree stmt; + + block = c_end_compound_stmt (loc, block, true); + + stmt = make_node (OACC_DATA); + TREE_TYPE (stmt) = void_type_node; + OACC_DATA_CLAUSES (stmt) = clauses; + OACC_DATA_BODY (stmt) = block; + SET_EXPR_LOCATION (stmt, loc); + + return add_stmt (stmt); +} + /* Like c_begin_compound_stmt, except force the retention of the BLOCK. */ tree @@ -11761,6 +11818,7 @@ handle_omp_array_sections (tree c) OMP_CLAUSE_SIZE (c) = size; if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) return false; + gcc_assert (OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_FORCE_DEVICEPTR); tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); OMP_CLAUSE_MAP_KIND (c2) = OMP_CLAUSE_MAP_POINTER; if (!c_mark_addressable (t)) @@ -11824,7 +11882,7 @@ c_find_omp_placeholder_r (tree *tp, int *, void *data) return NULL_TREE; } -/* For all elements of CLAUSES, validate them vs OpenMP constraints. +/* For all elements of CLAUSES, validate them against their constraints. Remove any elements from the list that are invalid. */ tree @@ -12184,7 +12242,9 @@ c_finish_omp_clauses (tree clauses) else if (!c_mark_addressable (t)) remove = true; else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP - && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER) + && (OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER + || (OMP_CLAUSE_MAP_KIND (c) + == OMP_CLAUSE_MAP_FORCE_DEVICEPTR))) && !lang_hooks.types.omp_mappable_type (TREE_TYPE (t))) { error_at (OMP_CLAUSE_LOCATION (c), @@ -12253,6 +12313,11 @@ c_finish_omp_clauses (tree clauses) case OMP_CLAUSE_TASKGROUP: case OMP_CLAUSE_PROC_BIND: case OMP_CLAUSE__CILK_FOR_COUNT_: + case OMP_CLAUSE_NUM_GANGS: + case OMP_CLAUSE_NUM_WORKERS: + case OMP_CLAUSE_VECTOR_LENGTH: + case OMP_CLAUSE_ASYNC: + case OMP_CLAUSE_WAIT: pc = &OMP_CLAUSE_CHAIN (c); continue;