This is the mail archive of the gcc-patches@gcc.gnu.org mailing list for the GCC project.
Index Nav: | [Date Index] [Subject Index] [Author Index] [Thread Index] | |
---|---|---|
Message Nav: | [Date Prev] [Date Next] | [Thread Prev] [Thread Next] |
Other format: | [Raw text] |
Hi! On 11/13/2014 09:04 AM, Jakub Jelinek wrote:
On Wed, Nov 05, 2014 at 03:39:44PM -0600, James Norris wrote:* 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 should be on the above line, no need to wrap too early.
Fixed.
@@ -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;Shouldn't this be PRAGMA_OMP_CLAUSE_HOST (PRAGMA_OACC_CLAUSE_HOST) instead? It is _HOST in the C++ patch, are there no C tests with that clause covering it?
The "host" clause is a synonym for the "self" clause. The initial C++ patch did not treat "host" as a synonym and has amended accordingly.
+ 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. */Like in C++ patch, please avoid the comment, file a PR instead, or just queue the work for GCC 6.
Will submit a PR.
+ /* 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));build_int_cst not aligned below expr_loc.
Fixed.
+ if (CAN_HAVE_LOCATION_P (c)) + SET_EXPR_LOCATION (c, expr_loc); + if (c == boolean_true_node) + { + warning_at (expr_loc, 0, + "%<num_gangs%> value must be positive");This would fit perfectly on one line.
Fixed.
+ 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);Again, as in C++ patch, please avoid this. Use gomp-constants.h, or some enum, or explain what the -1 is, but avoid TODO XXX: FIX.
Fixed.
+ else + { + t = c_fully_fold (t, false, NULL); + }Please avoid the {}s and reindent.
Fixed.
-/* 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));I wonder if it wouldn't be better to just put the new OpenACC routines into a new block of code, not stick it in between the OpenMP handling routines, because diff apparently lost track and I'm afraid so will svn blame/git blame and we'll lose history for the OpenMP changes.
There was a mistake in naming the function: c_parser_omp_clause_vector_length.
Once it was renamed to: c_parser_oacc_clause_vector_length, diff was able to keep track.
+ case PRAGMA_OMP_CLAUSE_VECTOR_LENGTH: + clauses = c_parser_omp_clause_vector_length (parser, clauses); + c_name = "vector_length"; + break;That is OpenACC clause, right? Shouldn't the routine be called c_parser_oacc_clause_vector_length?
Fixed.
+ case PRAGMA_OMP_CLAUSE_WAIT: + clauses = c_parser_oacc_clause_wait (parser, clauses);E.g. c_parser_oacc_clause_wait is.
Fixed.
+ clauses = c_parser_oacc_all_clauses (parser, OACC_DATA_CLAUSE_MASK, + "#pragma acc data");Too many spaces after =.
Fixed.
+/* OpenACC 2.0: + # pragma acc kernels oacc-kernels-clause[optseq] new-line + structured-block + + LOC is the location of the #pragma token.Again, what is LOC?
Fixed by removal of comment, along with other occurences.
+ clauses = c_parser_oacc_all_clauses (parser, OACC_KERNELS_CLAUSE_MASK, + p_name);See above.
Fixed.
+ c_parser_error (parser, enter + ? "expected %<data%> in %<#pragma acc enter data%>" + : "expected %<data%> 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");See the C++ patch review.
Fixed.
+ 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");Similarly (though, in this case C++ had unconditional acc enter data).
Fixed.
+ clauses = c_parser_oacc_all_clauses (parser, OACC_PARALLEL_CLAUSE_MASK, + p_name);See above.
Fixed.
+ if (find_omp_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE) + { + error_at (loc, + "%<#pragma acc update%> must contain at least one " + "%<device%> or %<host/self%> clause");There is no host/self clause, so you better write or %<host%> or %<self%> clause?
Fixed.
Otherwise LGTM.
Thank you for taking the time to review! OK to commit after middle end is accepted? Regards, Jim
Attachment:
ChangeLog2.txt
Description: Text document
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index 8a4fd39..b71ff4d 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -71,6 +71,7 @@ along with GCC; see the file COPYING3. If not see #include "plugin.h" #include "omp-low.h" #include "builtins.h" +#include "gomp-constants.h" /* Initialization routine for this file. */ @@ -1239,10 +1240,15 @@ static vec<tree, va_gc> *c_parser_expr_list (c_parser *, bool, bool, vec<tree, va_gc> **, location_t *, tree *, vec<location_t> *, 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 +4488,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 +4842,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: @@ -9560,6 +9597,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) { @@ -9762,7 +9818,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 @@ -9789,20 +9845,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_OACC_CLAUSE_ASYNC; break; case 'c': if (!strcmp ("collapse", p)) result = PRAGMA_OMP_CLAUSE_COLLAPSE; + else if (!strcmp ("copy", p)) + result = PRAGMA_OACC_CLAUSE_COPY; else if (!strcmp ("copyin", p)) result = PRAGMA_OMP_CLAUSE_COPYIN; + else if (!strcmp ("copyout", p)) + result = PRAGMA_OACC_CLAUSE_COPYOUT; else if (!strcmp ("copyprivate", p)) result = PRAGMA_OMP_CLAUSE_COPYPRIVATE; + else if (!strcmp ("create", p)) + result = PRAGMA_OACC_CLAUSE_CREATE; break; case 'd': - if (!strcmp ("depend", p)) + if (!strcmp ("delete", p)) + result = PRAGMA_OACC_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_OACC_CLAUSE_DEVICEPTR; else if (!strcmp ("dist_schedule", p)) result = PRAGMA_OMP_CLAUSE_DIST_SCHEDULE; break; @@ -9814,6 +9882,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_OACC_CLAUSE_SELF; + break; case 'i': if (!strcmp ("inbranch", p)) result = PRAGMA_OMP_CLAUSE_INBRANCH; @@ -9837,10 +9909,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_OACC_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_OACC_CLAUSE_NUM_WORKERS; else if (flag_cilkplus && !strcmp ("nomask", p)) result = PRAGMA_CILK_CLAUSE_NOMASK; break; @@ -9851,6 +9927,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_OACC_CLAUSE_PRESENT; + else if (!strcmp ("present_or_copy", p) + || !strcmp ("pcopy", p)) + result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY; + else if (!strcmp ("present_or_copyin", p) + || !strcmp ("pcopyin", p)) + result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN; + else if (!strcmp ("present_or_copyout", p) + || !strcmp ("pcopyout", p)) + result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT; + else if (!strcmp ("present_or_create", p) + || !strcmp ("pcreate", p)) + result = PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE; else if (!strcmp ("private", p)) result = PRAGMA_OMP_CLAUSE_PRIVATE; else if (!strcmp ("proc_bind", p)) @@ -9871,6 +9961,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_OACC_CLAUSE_SELF; break; case 't': if (!strcmp ("taskgroup", p)) @@ -9887,9 +9979,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_OACC_CLAUSE_VECTOR_LENGTH; + else if (flag_cilkplus && !strcmp ("vectorlength", p)) result = PRAGMA_CILK_CLAUSE_VECTORLENGTH; break; + case 'w': + if (!strcmp ("wait", p)) + result = PRAGMA_OACC_CLAUSE_WAIT; + break; } } @@ -9916,7 +10014,58 @@ 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<tree, va_gc> *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 list 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 @@ -10023,7 +10172,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, @@ -10040,7 +10189,112 @@ 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_OACC_CLAUSE_COPY: + kind = OMP_CLAUSE_MAP_FORCE_TOFROM; + break; + case PRAGMA_OMP_CLAUSE_COPYIN: + kind = OMP_CLAUSE_MAP_FORCE_TO; + break; + case PRAGMA_OACC_CLAUSE_COPYOUT: + kind = OMP_CLAUSE_MAP_FORCE_FROM; + break; + case PRAGMA_OACC_CLAUSE_CREATE: + kind = OMP_CLAUSE_MAP_FORCE_ALLOC; + break; + case PRAGMA_OACC_CLAUSE_DELETE: + kind = OMP_CLAUSE_MAP_FORCE_DEALLOC; + break; + case PRAGMA_OMP_CLAUSE_DEVICE: + kind = OMP_CLAUSE_MAP_FORCE_TO; + break; + case PRAGMA_OACC_CLAUSE_SELF: + kind = OMP_CLAUSE_MAP_FORCE_FROM; + break; + case PRAGMA_OACC_CLAUSE_PRESENT: + kind = OMP_CLAUSE_MAP_FORCE_PRESENT; + break; + case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY: + kind = OMP_CLAUSE_MAP_TOFROM; + break; + case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN: + kind = OMP_CLAUSE_MAP_TO; + break; + case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT: + kind = OMP_CLAUSE_MAP_FROM; + break; + case PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE: + kind = OMP_CLAUSE_MAP_ALLOC; + 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); + + 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 @@ -10183,7 +10437,7 @@ c_parser_omp_clause_final (c_parser *parser, tree list) return list; } -/* OpenMP 2.5: +/* OpenACC, OpenMP 2.5: if ( expression ) */ static tree @@ -10251,6 +10505,50 @@ 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, "%<num_gangs%> 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 ) */ @@ -10296,6 +10594,100 @@ 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, + "%<num_workers%> 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; + + t = build_int_cst (integer_type_node, GOMP_ASYNC_NOVAL); + + 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 */ @@ -10547,6 +10939,51 @@ c_parser_omp_clause_untied (c_parser *parser ATTRIBUTE_UNUSED, tree list) return c; } +/* OpenACC: + vector_length ( expression ) */ + +static tree +c_parser_oacc_clause_vector_length (c_parser *parser, tree 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); + + 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, + "%<vector_length%> 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 */ @@ -11083,6 +11520,140 @@ 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_OACC_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_OACC_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_OACC_CLAUSE_COPYOUT: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "copyout"; + break; + case PRAGMA_OACC_CLAUSE_CREATE: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "create"; + break; + case PRAGMA_OACC_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_OACC_CLAUSE_DEVICEPTR: + clauses = c_parser_oacc_data_clause_deviceptr (parser, clauses); + c_name = "deviceptr"; + break; + case PRAGMA_OMP_CLAUSE_IF: + clauses = c_parser_omp_clause_if (parser, clauses); + c_name = "if"; + break; + case PRAGMA_OACC_CLAUSE_NUM_GANGS: + clauses = c_parser_omp_clause_num_gangs (parser, clauses); + c_name = "num_gangs"; + break; + case PRAGMA_OACC_CLAUSE_NUM_WORKERS: + clauses = c_parser_omp_clause_num_workers (parser, clauses); + c_name = "num_workers"; + break; + case PRAGMA_OACC_CLAUSE_PRESENT: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "present"; + break; + case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "present_or_copy"; + break; + case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "present_or_copyin"; + break; + case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "present_or_copyout"; + break; + case PRAGMA_OACC_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_OACC_CLAUSE_SELF: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "self"; + break; + case PRAGMA_OACC_CLAUSE_VECTOR_LENGTH: + clauses = c_parser_oacc_clause_vector_length (parser, clauses); + c_name = "vector_length"; + break; + case PRAGMA_OACC_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. */ @@ -11313,7 +11884,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 @@ -11329,6 +11900,321 @@ c_parser_omp_structured_block (c_parser *parser) return pop_stmt_list (stmt); } +/* OpenACC 2.0: + # pragma acc cache (variable-list) new-line +*/ + +static tree +c_parser_oacc_cache (location_t loc ATTRIBUTE_UNUSED, c_parser *parser) +{ + c_parser_omp_var_list_parens (parser, OMP_NO_CLAUSE_CACHE, NULL); + c_parser_skip_to_pragma_eol (parser); + + return NULL_TREE; +} + +/* OpenACC 2.0: + # pragma acc data oacc-data-clause[optseq] new-line + structured-block +*/ + +#define OACC_DATA_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYIN) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_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 +*/ + +#define OACC_KERNELS_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYIN) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_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 +*/ + +#define OACC_ENTER_DATA_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYIN) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) + +#define OACC_EXIT_DATA_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DELETE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_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; + const char *p = ""; + + c_parser_consume_pragma (parser); + + if (c_parser_next_token_is (parser, CPP_NAME)) + { + p = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value); + c_parser_consume_token (parser); + } + + if (strcmp (p, "data") != 0) + { + error_at (loc, "expected %<data%> after %<#pragma acc enter %s%>", + enter ? "enter" : "exit" ); + parser->error = true; + c_parser_skip_to_pragma_eol (parser); + return; + } + + 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, + "%<#pragma acc %s data%> has no data movement clause", + enter ? "enter" : "exit"); + 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 +*/ + +#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 +*/ + +#define OACC_PARALLEL_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYIN) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_REDUCTION) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_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_OACC_CLAUSE_ASYNC) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SELF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_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: " + "%<device%>, %<host%> or %<self%> 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 +*/ + +#define OACC_WAIT_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_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 @@ -11805,10 +12691,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, @@ -12061,6 +12948,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 @@ -13650,6 +14538,29 @@ c_parser_omp_construct (c_parser *parser) switch (p_kind) { + case PRAGMA_OACC_CACHE: + strcpy (p_name, "#pragma acc"); + stmt = c_parser_oacc_cache (loc, parser); + break; + 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 338ef44..99b263b 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -11228,6 +11228,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 @@ -11759,6 +11816,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)) @@ -11822,7 +11880,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 @@ -12182,7 +12240,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), @@ -12251,6 +12311,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;
Index Nav: | [Date Index] [Subject Index] [Author Index] [Thread Index] | |
---|---|---|
Message Nav: | [Date Prev] [Date Next] | [Thread Prev] [Thread Next] |