Hi! On Tue, 05 May 2015 10:54:02 +0200, I wrote: > In follow-up messages, I'll be posting the separated parts (for easier > review) of a next set of OpenACC changes that we'd like to commit. > ChangeLog updates not yet written; will do that before commit, obviously. gcc/c-family/c-common.c | 3 +- gcc/c-family/c-common.h | 2 + gcc/c-family/c-omp.c | 105 ++ gcc/c-family/c-pragma.c | 4 + gcc/c-family/c-pragma.h | 14 +- gcc/c/c-parser.c | 1353 ++++++++++++---- gcc/c/c-tree.h | 3 +- gcc/c/c-typeck.c | 112 +- gcc/cp/cp-gimplify.c | 3 +- gcc/cp/cp-tree.h | 3 +- gcc/cp/parser.c | 1382 +++++++++++++---- gcc/cp/parser.h | 4 + gcc/cp/pt.c | 43 +- gcc/cp/semantics.c | 151 +- diff --git gcc/c-family/c-common.c gcc/c-family/c-common.c index 9797e17..d89b348 100644 --- gcc/c-family/c-common.c +++ gcc/c-family/c-common.c @@ -809,7 +809,7 @@ const struct attribute_spec c_common_attribute_table[] = handle_omp_declare_simd_attribute, false }, { "cilk simd function", 0, -1, true, false, false, handle_omp_declare_simd_attribute, false }, - { "omp declare target", 0, 0, true, false, false, + { "omp declare target", 0, -1, true, false, false, handle_omp_declare_target_attribute, false }, { "alloc_align", 1, 1, false, true, true, handle_alloc_align_attribute, false }, @@ -823,6 +823,7 @@ const struct attribute_spec c_common_attribute_table[] = handle_bnd_legacy, false }, { "bnd_instrument", 0, 0, true, false, false, handle_bnd_instrument, false }, + { "oacc declare", 0, -1, true, false, false, NULL, false }, { NULL, 0, 0, false, false, false, NULL, false } }; diff --git gcc/c-family/c-common.h gcc/c-family/c-common.h index 603d3f0..fcaebca 100644 --- gcc/c-family/c-common.h +++ gcc/c-family/c-common.h @@ -1249,6 +1249,8 @@ extern void c_omp_split_clauses (location_t, enum tree_code, omp_clause_mask, extern tree c_omp_declare_simd_clauses_to_numbers (tree, tree); extern void c_omp_declare_simd_clauses_to_decls (tree, tree); extern enum omp_clause_default_kind c_omp_predetermined_sharing (tree); +extern int oacc_extract_device_id (const char *); +extern tree oacc_filter_device_types (tree); /* Return next tree in the chain for chain_next walking of tree nodes. */ static inline tree diff --git gcc/c-family/c-omp.c gcc/c-family/c-omp.c index 86a9f54..1c82bf5 100644 --- gcc/c-family/c-omp.c +++ gcc/c-family/c-omp.c @@ -1087,3 +1087,108 @@ c_omp_predetermined_sharing (tree decl) return OMP_CLAUSE_DEFAULT_UNSPECIFIED; } + +/* Return a numerical code representing the device_type. Currently, + only device_type(nvidia) is supported. All device_type parameters + are treated as case-insensitive keywords. */ + +int +oacc_extract_device_id (const char *device) +{ + if (!strcasecmp (device, "nvidia")) + return GOMP_DEVICE_NVIDIA_PTX; + return GOMP_DEVICE_NONE; +} + +/* Filter out the list of unsupported OpenACC device_types. */ + +tree +oacc_filter_device_types (tree clauses) +{ + tree c, prev; + tree dtype = NULL_TREE; + tree seen_nvidia = NULL_TREE; + tree seen_default = NULL_TREE; + + /* First scan for all device_type clauses. */ + for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) + { + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEVICE_TYPE) + { + int code = TREE_INT_CST_LOW (OMP_CLAUSE_DEVICE_TYPE_DEVICES (c)); + + if (code == GOMP_DEVICE_DEFAULT) + { + if (seen_default) + { + seen_default = NULL_TREE; + error_at (OMP_CLAUSE_LOCATION (c), + "duplicate device_type (*)"); + goto filter_error; + } + else + seen_default = OMP_CLAUSE_DEVICE_TYPE_CLAUSES (c); + } + if (code & (1 << GOMP_DEVICE_NVIDIA_PTX)) + { + if (seen_nvidia) + { + seen_nvidia = NULL_TREE; + error_at (OMP_CLAUSE_LOCATION (c), + "duplicate device_type (nvidia)"); + goto filter_error; + } + else + seen_nvidia = OMP_CLAUSE_DEVICE_TYPE_CLAUSES (c); + } + } + } + + /* Don't do anything if there aren't any device_type clauses. */ + if (seen_nvidia == NULL_TREE && seen_default == NULL_TREE) + return clauses; + + dtype = seen_nvidia ? seen_nvidia : seen_default; + + /* Now filter out clauses if necessary. */ + for (c = clauses; c && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEVICE_TYPE; + c = OMP_CLAUSE_CHAIN (c)) + { + tree t; + + prev = NULL_TREE; + + for (t = dtype; t; t = OMP_CLAUSE_CHAIN (t)) + { + if (OMP_CLAUSE_CODE (t) == OMP_CLAUSE_CODE (c)) + { + /* Remove c from clauses. */ + tree next = OMP_CLAUSE_CHAIN (c); + + if (prev) + OMP_CLAUSE_CHAIN (prev) = next; + + break; + } + } + + prev = c; + } + + filter_error: + /* Remove all device_type clauses. Those clauses are located at the + beginning of the clause list. */ + for (c = clauses; c && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEVICE_TYPE; + c = OMP_CLAUSE_CHAIN (c)) + ; + + if (c == NULL_TREE) + return dtype; + + clauses = c; + for (prev = c, c = OMP_CLAUSE_CHAIN (c); c; c = OMP_CLAUSE_CHAIN (c)) + prev = c; + + OMP_CLAUSE_CHAIN (prev) = dtype; + return clauses; +} diff --git gcc/c-family/c-pragma.c gcc/c-family/c-pragma.c index 6894f0e..a1e8da3 100644 --- gcc/c-family/c-pragma.c +++ gcc/c-family/c-pragma.c @@ -1194,13 +1194,17 @@ static vec registered_pp_pragmas; struct omp_pragma_def { const char *name; unsigned int id; }; static const struct omp_pragma_def oacc_pragmas[] = { + { "atomic", PRAGMA_OACC_ATOMIC }, { "cache", PRAGMA_OACC_CACHE }, { "data", PRAGMA_OACC_DATA }, + { "declare", PRAGMA_OACC_DECLARE }, { "enter", PRAGMA_OACC_ENTER_DATA }, { "exit", PRAGMA_OACC_EXIT_DATA }, + { "host_data", PRAGMA_OACC_HOST_DATA }, { "kernels", PRAGMA_OACC_KERNELS }, { "loop", PRAGMA_OACC_LOOP }, { "parallel", PRAGMA_OACC_PARALLEL }, + { "routine", PRAGMA_OACC_ROUTINE }, { "update", PRAGMA_OACC_UPDATE }, { "wait", PRAGMA_OACC_WAIT } }; diff --git gcc/c-family/c-pragma.h gcc/c-family/c-pragma.h index eff94c1..fe4c168 100644 --- gcc/c-family/c-pragma.h +++ gcc/c-family/c-pragma.h @@ -27,13 +27,17 @@ along with GCC; see the file COPYING3. If not see typedef enum pragma_kind { PRAGMA_NONE = 0, + PRAGMA_OACC_ATOMIC, PRAGMA_OACC_CACHE, PRAGMA_OACC_DATA, + PRAGMA_OACC_DECLARE, PRAGMA_OACC_ENTER_DATA, PRAGMA_OACC_EXIT_DATA, + PRAGMA_OACC_HOST_DATA, PRAGMA_OACC_KERNELS, PRAGMA_OACC_LOOP, PRAGMA_OACC_PARALLEL, + PRAGMA_OACC_ROUTINE, PRAGMA_OACC_UPDATE, PRAGMA_OACC_WAIT, PRAGMA_OMP_ATOMIC, @@ -132,13 +136,19 @@ typedef enum pragma_omp_clause { /* Clauses for OpenACC. */ PRAGMA_OACC_CLAUSE_ASYNC = PRAGMA_CILK_CLAUSE_VECTORLENGTH + 1, PRAGMA_OACC_CLAUSE_AUTO, + PRAGMA_OACC_CLAUSE_BIND, PRAGMA_OACC_CLAUSE_COPY, PRAGMA_OACC_CLAUSE_COPYOUT, PRAGMA_OACC_CLAUSE_CREATE, PRAGMA_OACC_CLAUSE_DELETE, PRAGMA_OACC_CLAUSE_DEVICEPTR, + PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT, + PRAGMA_OACC_CLAUSE_DEVICE_TYPE, PRAGMA_OACC_CLAUSE_GANG, PRAGMA_OACC_CLAUSE_HOST, + PRAGMA_OACC_CLAUSE_INDEPENDENT, + PRAGMA_OACC_CLAUSE_LINK, + PRAGMA_OACC_CLAUSE_NOHOST, PRAGMA_OACC_CLAUSE_NUM_GANGS, PRAGMA_OACC_CLAUSE_NUM_WORKERS, PRAGMA_OACC_CLAUSE_PRESENT, @@ -146,8 +156,9 @@ typedef enum pragma_omp_clause { PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN, PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT, PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE, - PRAGMA_OACC_CLAUSE_SELF, PRAGMA_OACC_CLAUSE_SEQ, + PRAGMA_OACC_CLAUSE_TILE, + PRAGMA_OACC_CLAUSE_USE_DEVICE, PRAGMA_OACC_CLAUSE_VECTOR, PRAGMA_OACC_CLAUSE_VECTOR_LENGTH, PRAGMA_OACC_CLAUSE_WAIT, @@ -155,6 +166,7 @@ typedef enum pragma_omp_clause { PRAGMA_OACC_CLAUSE_COLLAPSE = PRAGMA_OMP_CLAUSE_COLLAPSE, PRAGMA_OACC_CLAUSE_COPYIN = PRAGMA_OMP_CLAUSE_COPYIN, PRAGMA_OACC_CLAUSE_DEVICE = PRAGMA_OMP_CLAUSE_DEVICE, + PRAGMA_OACC_CLAUSE_DEFAULT = PRAGMA_OMP_CLAUSE_DEFAULT, PRAGMA_OACC_CLAUSE_FIRSTPRIVATE = PRAGMA_OMP_CLAUSE_FIRSTPRIVATE, PRAGMA_OACC_CLAUSE_IF = PRAGMA_OMP_CLAUSE_IF, PRAGMA_OACC_CLAUSE_PRIVATE = PRAGMA_OMP_CLAUSE_PRIVATE, diff --git gcc/c/c-parser.c gcc/c/c-parser.c index 015de7f..a1543a7 100644 --- gcc/c/c-parser.c +++ gcc/c/c-parser.c @@ -240,6 +240,10 @@ typedef struct GTY(()) c_parser { /* Buffer to hold all the tokens from parsing the vector attribute for the SIMD-enabled functions (formerly known as elemental functions). */ vec *cilk_simd_fn_tokens; + + /* OpenACC specific parser information. */ + + vec *oacc_routines; } c_parser; @@ -1181,7 +1185,8 @@ enum c_parser_prec { static void c_parser_external_declaration (c_parser *); static void c_parser_asm_definition (c_parser *); static void c_parser_declaration_or_fndef (c_parser *, bool, bool, bool, - bool, bool, tree *, vec); + bool, bool, tree *, vec, + tree, bool); static void c_parser_static_assert_declaration_no_semi (c_parser *); static void c_parser_static_assert_declaration (c_parser *); static void c_parser_declspecs (c_parser *, struct c_declspecs *, bool, bool, @@ -1252,7 +1257,8 @@ static vec *c_parser_expr_list (c_parser *, bool, bool, unsigned int * = NULL); static void c_parser_oacc_enter_exit_data (c_parser *, bool); static void c_parser_oacc_update (c_parser *); -static tree c_parser_oacc_loop (location_t, c_parser *, char *); +static tree c_parser_oacc_loop (location_t, c_parser *, char *, + omp_clause_mask, tree *); static void c_parser_omp_construct (c_parser *); static void c_parser_omp_threadprivate (c_parser *); static void c_parser_omp_barrier (c_parser *); @@ -1270,6 +1276,9 @@ static bool c_parser_pragma (c_parser *, enum pragma_context); static bool c_parser_omp_target (c_parser *, enum pragma_context); static void c_parser_omp_end_declare_target (c_parser *); static void c_parser_omp_declare (c_parser *, enum pragma_context); +static void c_parser_oacc_routine (c_parser *parser, enum pragma_context + context); +static void c_parser_oacc_declare (c_parser *parser); /* These Objective-C parser functions are only ever called when compiling Objective-C. */ @@ -1306,6 +1315,11 @@ static tree c_parser_array_notation (location_t, c_parser *, tree, tree); static tree c_parser_cilk_clause_vectorlength (c_parser *, tree, bool); static void c_parser_cilk_grainsize (c_parser *); +/* OpenACC support. */ +static tree c_parser_oacc_all_clauses (c_parser *, omp_clause_mask, + const char *, omp_clause_mask, + bool, bool); + /* Parse a translation unit (C90 6.7, C99 6.9). translation-unit: @@ -1449,12 +1463,13 @@ c_parser_external_declaration (c_parser *parser) only tell which after parsing the declaration specifiers, if any, and the first declarator. */ c_parser_declaration_or_fndef (parser, true, true, true, false, true, - NULL, vNULL); + NULL, vNULL, NULL_TREE, false); break; } } static void c_finish_omp_declare_simd (c_parser *, tree, tree, vec); +static void c_finish_oacc_routine (c_parser *, tree, tree, bool); /* Parse a declaration or function definition (C90 6.5, 6.7.1, C99 6.7, 6.9.1). If FNDEF_OK is true, a function definition is @@ -1532,7 +1547,8 @@ c_parser_declaration_or_fndef (c_parser *parser, bool fndef_ok, bool static_assert_ok, bool empty_ok, bool nested, bool start_attr_ok, tree *objc_foreach_object_declaration, - vec omp_declare_simd_clauses) + vec omp_declare_simd_clauses, + tree oacc_routine_clauses, bool oacc_routine_named) { struct c_declspecs *specs; tree prefix_attrs; @@ -1710,6 +1726,9 @@ c_parser_declaration_or_fndef (c_parser *parser, bool fndef_ok, || !vec_safe_is_empty (parser->cilk_simd_fn_tokens)) c_finish_omp_declare_simd (parser, NULL_TREE, NULL_TREE, omp_declare_simd_clauses); + else + c_finish_oacc_routine (parser, NULL_TREE, + oacc_routine_clauses, oacc_routine_named); c_parser_skip_to_end_of_block_or_statement (parser); return; } @@ -1806,6 +1825,9 @@ c_parser_declaration_or_fndef (c_parser *parser, bool fndef_ok, || !vec_safe_is_empty (parser->cilk_simd_fn_tokens)) c_finish_omp_declare_simd (parser, d, NULL_TREE, omp_declare_simd_clauses); + else + c_finish_oacc_routine (parser, d, oacc_routine_clauses, + oacc_routine_named); } else { @@ -1819,6 +1841,10 @@ c_parser_declaration_or_fndef (c_parser *parser, bool fndef_ok, || !vec_safe_is_empty (parser->cilk_simd_fn_tokens)) c_finish_omp_declare_simd (parser, d, NULL_TREE, omp_declare_simd_clauses); + else + c_finish_oacc_routine (parser, d, oacc_routine_clauses, + oacc_routine_named); + start_init (d, asm_name, global_bindings_p ()); init_loc = c_parser_peek_token (parser)->location; init = c_parser_initializer (parser); @@ -1864,6 +1890,9 @@ c_parser_declaration_or_fndef (c_parser *parser, bool fndef_ok, temp_store_parm_decls (d, parms); c_finish_omp_declare_simd (parser, d, parms, omp_declare_simd_clauses); + c_finish_oacc_routine (parser, d, oacc_routine_clauses, + oacc_routine_named); + if (parms) temp_pop_parm_decls (); } @@ -1970,13 +1999,17 @@ c_parser_declaration_or_fndef (c_parser *parser, bool fndef_ok, function definitions either. */ while (c_parser_next_token_is_not (parser, CPP_EOF) && c_parser_next_token_is_not (parser, CPP_OPEN_BRACE)) - c_parser_declaration_or_fndef (parser, false, false, false, - true, false, NULL, vNULL); + c_parser_declaration_or_fndef (parser, false, false, false, true, + false, NULL, vNULL, NULL_TREE, false); store_parm_decls (); if (omp_declare_simd_clauses.exists () || !vec_safe_is_empty (parser->cilk_simd_fn_tokens)) c_finish_omp_declare_simd (parser, current_function_decl, NULL_TREE, omp_declare_simd_clauses); + else + c_finish_oacc_routine (parser, current_function_decl, + oacc_routine_clauses, oacc_routine_named); + DECL_STRUCT_FUNCTION (current_function_decl)->function_start_locus = c_parser_peek_token (parser)->location; fnbody = c_parser_compound_statement (parser); @@ -4624,7 +4657,7 @@ c_parser_compound_statement_nostart (c_parser *parser) last_label = false; mark_valid_location_for_stdc_pragma (false); c_parser_declaration_or_fndef (parser, true, true, true, true, - true, NULL, vNULL); + true, NULL, vNULL, NULL_TREE, false); if (last_stmt) pedwarn_c90 (loc, OPT_Wdeclaration_after_statement, "ISO C90 forbids mixed declarations and code"); @@ -4649,7 +4682,8 @@ c_parser_compound_statement_nostart (c_parser *parser) last_label = false; mark_valid_location_for_stdc_pragma (false); c_parser_declaration_or_fndef (parser, true, true, true, true, - true, NULL, vNULL); + true, NULL, vNULL, NULL_TREE, + false); /* Following the old parser, __extension__ does not disable this diagnostic. */ restore_extension_diagnostics (ext); @@ -4798,7 +4832,7 @@ c_parser_label (c_parser *parser) /*static_assert_ok*/ true, /*empty_ok*/ true, /*nested*/ true, /*start_attr_ok*/ true, NULL, - vNULL); + vNULL, NULL_TREE, false); } } } @@ -5501,7 +5535,8 @@ c_parser_for_statement (c_parser *parser, bool ivdep) else if (c_parser_next_tokens_start_declaration (parser)) { c_parser_declaration_or_fndef (parser, true, true, true, true, true, - &object_expression, vNULL); + &object_expression, vNULL, NULL_TREE, + false); parser->objc_could_be_foreach_context = false; if (c_parser_next_token_is_keyword (parser, RID_IN)) @@ -5530,7 +5565,8 @@ c_parser_for_statement (c_parser *parser, bool ivdep) ext = disable_extension_diagnostics (); c_parser_consume_token (parser); c_parser_declaration_or_fndef (parser, true, true, true, true, - true, &object_expression, vNULL); + true, &object_expression, vNULL, + NULL_TREE, false); parser->objc_could_be_foreach_context = false; restore_extension_diagnostics (ext); @@ -8658,8 +8694,9 @@ c_parser_objc_methodprotolist (c_parser *parser) c_parser_consume_token (parser); } else - c_parser_declaration_or_fndef (parser, false, false, true, - false, true, NULL, vNULL); + c_parser_declaration_or_fndef (parser, false, false, true,false, + true, NULL, vNULL, NULL_TREE, + false); break; } } @@ -9608,14 +9645,36 @@ c_parser_pragma (c_parser *parser, enum pragma_context context) switch (id) { + case PRAGMA_OACC_DECLARE: + c_parser_oacc_declare (parser); + return false; + case PRAGMA_OACC_ENTER_DATA: + if (context != pragma_compound) + { + if (context == pragma_stmt) + c_parser_error (parser, "%<#pragma acc enter data%> may only be " + "used in compound statements"); + goto bad_stmt; + } c_parser_oacc_enter_exit_data (parser, true); return false; case PRAGMA_OACC_EXIT_DATA: + if (context != pragma_compound) + { + if (context == pragma_stmt) + c_parser_error (parser, "%<#pragma acc exit data%> may only be " + "used in compound statements"); + goto bad_stmt; + } c_parser_oacc_enter_exit_data (parser, false); return false; + case PRAGMA_OACC_ROUTINE: + c_parser_oacc_routine (parser, context); + return false; + case PRAGMA_OACC_UPDATE: if (context != pragma_compound) { @@ -9761,6 +9820,16 @@ c_parser_pragma (c_parser *parser, enum pragma_context context) c_parser_cilk_grainsize (parser); return false; + case PRAGMA_OACC_WAIT: + if (context != pragma_compound) + { + if (context == pragma_stmt) + c_parser_error (parser, "%<#pragma acc enter data%> may only be " + "used in compound statements"); + goto bad_stmt; + } + /* FALL THROUGH. */ + default: if (id < PRAGMA_FIRST_EXTERNAL) { @@ -9837,7 +9906,7 @@ c_parser_pragma_pch_preprocess (c_parser *parser) returned and the token is consumed. */ static pragma_omp_clause -c_parser_omp_clause_name (c_parser *parser) +c_parser_omp_clause_name (c_parser *parser, bool consume_token = true) { pragma_omp_clause result = PRAGMA_OMP_CLAUSE_NONE; @@ -9861,6 +9930,10 @@ c_parser_omp_clause_name (c_parser *parser) else if (!strcmp ("async", p)) result = PRAGMA_OACC_CLAUSE_ASYNC; break; + case 'b': + if (!strcmp ("bind", p)) + result = PRAGMA_OACC_CLAUSE_BIND; + break; case 'c': if (!strcmp ("collapse", p)) result = PRAGMA_OMP_CLAUSE_COLLAPSE; @@ -9882,6 +9955,11 @@ c_parser_omp_clause_name (c_parser *parser) result = PRAGMA_OMP_CLAUSE_DEPEND; else if (!strcmp ("device", p)) result = PRAGMA_OMP_CLAUSE_DEVICE; + else if (!strcmp ("device_resident", p)) + result = PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT; + else if (!strcmp ("device_type", p) + || !strcmp ("dtype", p)) + result = PRAGMA_OACC_CLAUSE_DEVICE_TYPE; else if (!strcmp ("deviceptr", p)) result = PRAGMA_OACC_CLAUSE_DEVICEPTR; else if (!strcmp ("dist_schedule", p)) @@ -9906,12 +9984,16 @@ c_parser_omp_clause_name (c_parser *parser) case 'i': if (!strcmp ("inbranch", p)) result = PRAGMA_OMP_CLAUSE_INBRANCH; + else if (!strcmp ("independent", p)) + result = PRAGMA_OACC_CLAUSE_INDEPENDENT; break; case 'l': if (!strcmp ("lastprivate", p)) result = PRAGMA_OMP_CLAUSE_LASTPRIVATE; else if (!strcmp ("linear", p)) result = PRAGMA_OMP_CLAUSE_LINEAR; + else if (!strcmp ("link", p)) + result = PRAGMA_OACC_CLAUSE_LINK; break; case 'm': if (!strcmp ("map", p)) @@ -9926,6 +10008,8 @@ 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 ("nohost", p)) + result = PRAGMA_OACC_CLAUSE_NOHOST; else if (!strcmp ("num_gangs", p)) result = PRAGMA_OACC_CLAUSE_NUM_GANGS; else if (!strcmp ("num_teams", p)) @@ -9974,20 +10058,22 @@ c_parser_omp_clause_name (c_parser *parser) result = PRAGMA_OMP_CLAUSE_SCHEDULE; else if (!strcmp ("sections", p)) result = PRAGMA_OMP_CLAUSE_SECTIONS; + else if (!strcmp ("self", p)) /* "self" is a synonym for "host". */ + result = PRAGMA_OACC_CLAUSE_HOST; else if (!strcmp ("seq", p)) result = PRAGMA_OACC_CLAUSE_SEQ; else if (!strcmp ("shared", p)) 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)) result = PRAGMA_OMP_CLAUSE_TASKGROUP; else if (!strcmp ("thread_limit", p)) result = PRAGMA_OMP_CLAUSE_THREAD_LIMIT; + else if (!strcmp ("tile", p)) + result = PRAGMA_OACC_CLAUSE_TILE; else if (!strcmp ("to", p)) result = PRAGMA_OMP_CLAUSE_TO; break; @@ -9996,6 +10082,8 @@ c_parser_omp_clause_name (c_parser *parser) result = PRAGMA_OMP_CLAUSE_UNIFORM; else if (!strcmp ("untied", p)) result = PRAGMA_OMP_CLAUSE_UNTIED; + else if (!strcmp ("use_device", p)) + result = PRAGMA_OACC_CLAUSE_USE_DEVICE; break; case 'v': if (!strcmp ("vector", p)) @@ -10014,7 +10102,7 @@ c_parser_omp_clause_name (c_parser *parser) } } - if (result != PRAGMA_OMP_CLAUSE_NONE) + if (consume_token && result != PRAGMA_OMP_CLAUSE_NONE) c_parser_consume_token (parser); return result; @@ -10053,7 +10141,8 @@ c_parser_oacc_wait_list (c_parser *parser, location_t clause_loc, tree list) if (args->length () == 0) { - c_parser_error (parser, "expected integer expression before ')'"); + c_parser_error (parser, + "expected integer expression list before %<)%>"); release_tree_vector (args); return list; } @@ -10245,6 +10334,8 @@ c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind, copyout ( variable-list ) create ( variable-list ) delete ( variable-list ) + device_resident ( variable-list ) + link ( variable-list ) present ( variable-list ) present_or_copy ( variable-list ) pcopy ( variable-list ) @@ -10280,10 +10371,15 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind, case PRAGMA_OACC_CLAUSE_DEVICE: kind = GOMP_MAP_FORCE_TO; break; + case PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT: + kind = GOMP_MAP_DEVICE_RESIDENT; + break; case PRAGMA_OACC_CLAUSE_HOST: - case PRAGMA_OACC_CLAUSE_SELF: kind = GOMP_MAP_FORCE_FROM; break; + case PRAGMA_OACC_CLAUSE_LINK: + kind = GOMP_MAP_LINK; + break; case PRAGMA_OACC_CLAUSE_PRESENT: kind = GOMP_MAP_FORCE_PRESENT; break; @@ -10410,7 +10506,8 @@ c_parser_omp_clause_copyprivate (c_parser *parser, tree list) default ( shared | none ) */ static tree -c_parser_omp_clause_default (c_parser *parser, tree list) +c_parser_omp_clause_default (c_parser *parser, tree list, + bool only_none = false) { enum omp_clause_default_kind kind = OMP_CLAUSE_DEFAULT_UNSPECIFIED; location_t loc = c_parser_peek_token (parser)->location; @@ -10431,7 +10528,7 @@ c_parser_omp_clause_default (c_parser *parser, tree list) break; case 's': - if (strcmp ("shared", p) != 0) + if (strcmp ("shared", p) != 0 || only_none) goto invalid_kind; kind = OMP_CLAUSE_DEFAULT_SHARED; break; @@ -10445,7 +10542,10 @@ c_parser_omp_clause_default (c_parser *parser, tree list) else { invalid_kind: - c_parser_error (parser, "expected % or %"); + if (only_none) + c_parser_error (parser, "expected %"); + else + c_parser_error (parser, "expected % or %"); } c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, "expected %<)%>"); @@ -10562,139 +10662,195 @@ c_parser_omp_clause_nowait (c_parser *parser ATTRIBUTE_UNUSED, tree list) return c; } -/* OpenACC: - num_gangs ( expression ) */ +/* Attempt to statically determine when the number T isn't positive. + Warn if we determined this and return positive one as the new + expression. */ static tree -c_parser_omp_clause_num_gangs (c_parser *parser, tree list) +require_positive_expr (tree t, location_t loc, const char *str) { - location_t num_gangs_loc = c_parser_peek_token (parser)->location; - if (c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>")) + tree c = fold_build2_loc (loc, LE_EXPR, boolean_type_node, t, + build_int_cst (TREE_TYPE (t), 0)); + if (c == boolean_true_node) { - 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; + warning_at (loc, 0, + "%<%s%> value must be positive", str); + t = integer_one_node; } - - return list; + return t; } -/* OpenMP 2.5: +/* OpenACC: + num_gangs ( expression ) + num_workers ( expression ) + vector_length ( expression ) + + OpenMP 2.5: num_threads ( expression ) */ static tree -c_parser_omp_clause_num_threads (c_parser *parser, tree list) +c_parser_omp_positive_int_clause (c_parser *parser, pragma_omp_clause c_kind, + const char *str, tree list) { - location_t num_threads_loc = c_parser_peek_token (parser)->location; - if (c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>")) + omp_clause_code kind; + switch (c_kind) { - 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 %<)%>"); + default: + gcc_unreachable (); + case PRAGMA_OACC_CLAUSE_NUM_GANGS: + kind = OMP_CLAUSE_NUM_GANGS; + break; + case PRAGMA_OMP_CLAUSE_NUM_THREADS: + kind = OMP_CLAUSE_NUM_THREADS; + break; + case PRAGMA_OACC_CLAUSE_NUM_WORKERS: + kind = OMP_CLAUSE_NUM_WORKERS; + break; + case PRAGMA_OACC_CLAUSE_VECTOR_LENGTH: + kind = OMP_CLAUSE_VECTOR_LENGTH; + break; + } - if (!INTEGRAL_TYPE_P (TREE_TYPE (t))) - { - c_parser_error (parser, "expected integer expression"); - return list; - } + location_t loc = c_parser_peek_token (parser)->location; + if (!c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>")) + 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; - } + 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); - check_no_duplicate_clause (list, OMP_CLAUSE_NUM_THREADS, "num_threads"); + c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, "expected %<)%>"); - c = build_omp_clause (num_threads_loc, OMP_CLAUSE_NUM_THREADS); - OMP_CLAUSE_NUM_THREADS_EXPR (c) = t; - OMP_CLAUSE_CHAIN (c) = list; - list = c; + if (!INTEGRAL_TYPE_P (TREE_TYPE (t))) + { + c_parser_error (parser, "expected integer expression"); + return list; } - return list; + require_positive_expr (t, expr_loc, str); + + check_no_duplicate_clause (list, kind, str); + + c = build_omp_clause (loc, kind); + OMP_CLAUSE_OPERAND (c, 0) = t; + OMP_CLAUSE_CHAIN (c) = list; + return c; } /* OpenACC: - num_workers ( expression ) */ + gang [( gang_expr_list )] + worker [( expression )] + vector [( expression )] */ static tree -c_parser_omp_clause_num_workers (c_parser *parser, tree list) +c_parser_oacc_shape_clause (c_parser *parser, pragma_omp_clause c_kind, + const char *str, tree list) { - location_t num_workers_loc = c_parser_peek_token (parser)->location; - if (c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>")) + omp_clause_code kind; + const char *id = "num"; + + switch (c_kind) { - 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); + default: + gcc_unreachable (); + case PRAGMA_OACC_CLAUSE_GANG: + kind = OMP_CLAUSE_GANG; + break; + case PRAGMA_OACC_CLAUSE_VECTOR: + kind = OMP_CLAUSE_VECTOR; + id = "length"; + break; + case PRAGMA_OACC_CLAUSE_WORKER: + kind = OMP_CLAUSE_WORKER; + break; + } - c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, "expected %<)%>"); + tree op0 = NULL_TREE, op1 = NULL_TREE; + location_t loc = c_parser_peek_token (parser)->location; - if (!INTEGRAL_TYPE_P (TREE_TYPE (t))) - { - c_parser_error (parser, "expected integer expression"); - return list; - } + if (c_parser_next_token_is (parser, CPP_OPEN_PAREN)) + { + tree *op_to_parse = &op0; + c_parser_consume_token (parser); - /* 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) + do { - warning_at (expr_loc, 0, - "% value must be positive"); - t = integer_one_node; - } + if (c_parser_next_token_is (parser, CPP_NAME) + || c_parser_next_token_is (parser, CPP_KEYWORD)) + { + tree name_kind = c_parser_peek_token (parser)->value; + const char *p = IDENTIFIER_POINTER (name_kind); + if (kind == OMP_CLAUSE_GANG && strcmp ("static", p) == 0) + { + c_parser_consume_token (parser); + if (!c_parser_require (parser, CPP_COLON, "expected %<:%>")) + { + c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, 0); + return list; + } + op_to_parse = &op1; + if (c_parser_next_token_is (parser, CPP_MULT)) + { + c_parser_consume_token (parser); + *op_to_parse = integer_minus_one_node; + continue; + } + } + else if (strcmp (id, p) == 0) + { + c_parser_consume_token (parser); + if (!c_parser_require (parser, CPP_COLON, "expected %<:%>")) + { + c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, 0); + return list; + } + } + else + { + if (kind == OMP_CLAUSE_GANG) + c_parser_error (parser, "expected %<%num%> or %"); + else if (kind == OMP_CLAUSE_VECTOR) + c_parser_error (parser, "expected %"); + else + c_parser_error (parser, "expected %"); + c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, 0); + return list; + } + } + + if (*op_to_parse != NULL_TREE) + { + c_parser_error (parser, "duplicate operand to clause"); + c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, 0); + return list; + } - check_no_duplicate_clause (list, OMP_CLAUSE_NUM_WORKERS, "num_workers"); + location_t expr_loc = c_parser_peek_token (parser)->location; + tree expr = c_parser_expression (parser).value; + if (expr == error_mark_node) + { + c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, 0); + return list; + } - 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; + mark_exp_read (expr); + require_positive_expr (expr, expr_loc, str); + *op_to_parse = expr; + } + while (!c_parser_next_token_is (parser, CPP_CLOSE_PAREN)); + c_parser_consume_token (parser); } - return list; + check_no_duplicate_clause (list, kind, str); + + tree c = build_omp_clause (loc, kind); + if (op0) + OMP_CLAUSE_OPERAND (c, 0) = op0; + if (op1) + OMP_CLAUSE_OPERAND (c, 1) = op1; + OMP_CLAUSE_CHAIN (c) = list; + return c; } /* OpenACC: @@ -10732,6 +10888,195 @@ c_parser_oacc_clause_async (c_parser *parser, tree list) return list; } +/* OpenACC 2.0: + bind ( identifier ) + bind ( string-literal ) */ + +static tree +c_parser_oacc_clause_bind (c_parser *parser, tree list) +{ + location_t loc = c_parser_peek_token (parser)->location; + + parser->lex_untranslated_string = true; + if (!c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>")) + { + parser->lex_untranslated_string = false; + return list; + } + if (c_parser_next_token_is (parser, CPP_NAME) + || c_parser_next_token_is (parser, CPP_STRING)) + { + tree t = c_parser_peek_token (parser)->value; + c_parser_consume_token (parser); + tree c = build_omp_clause (loc, OMP_CLAUSE_BIND); + OMP_CLAUSE_BIND_NAME (c) = t; + OMP_CLAUSE_CHAIN (c) = list; + list = c; + } + else + c_parser_error (parser, "expected identifier or character string literal"); + parser->lex_untranslated_string = false; + c_parser_require (parser, CPP_CLOSE_PAREN, "expected %<)%>"); + return list; +} + +/* OpenACC 2.0: + device_type ( size-expr-list ) clauses */ + +static tree +c_parser_oacc_clause_device_type (c_parser *parser, omp_clause_mask mask, + tree list) +{ + tree c, clauses; + location_t loc; + int dev_id = GOMP_DEVICE_NONE; + + loc = c_parser_peek_token (parser)->location; + if (!c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>")) + return list; + + if (c_parser_next_token_is (parser, CPP_MULT)) + { + c_parser_consume_token (parser); + dev_id = GOMP_DEVICE_DEFAULT; + if (!c_parser_require (parser, CPP_CLOSE_PAREN, "expected %<)%>")) + return list; + } + else + { + do + { + tree keyword = error_mark_node; + int dev = 0; + + if (c_parser_next_token_is (parser, CPP_NAME)) + { + keyword = c_parser_peek_token (parser)->value; + c_parser_consume_token (parser); + } + + if (keyword == error_mark_node) + { + error_at (loc, "expected keyword or %<)%>"); + c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, + "expected %<)%>"); + return list; + } + + dev = oacc_extract_device_id (IDENTIFIER_POINTER (keyword)); + if (dev) + dev_id |= 1 << dev; + + if (c_parser_next_token_is (parser, CPP_COMMA)) + c_parser_consume_token (parser); + } + while (c_parser_next_token_is_not (parser, CPP_CLOSE_PAREN)); + + /* Consume the trailing ')'. */ + c_parser_consume_token (parser); + } + + c = build_omp_clause (loc, OMP_CLAUSE_DEVICE_TYPE); + clauses = c_parser_oacc_all_clauses (parser, mask, "device_type", 0, false, + false); + OMP_CLAUSE_DEVICE_TYPE_CLAUSES (c) = clauses; + OMP_CLAUSE_DEVICE_TYPE_DEVICES (c) = build_int_cst (integer_type_node, + dev_id); + OMP_CLAUSE_CHAIN (c) = list; + return c; +} + +/* OpenACC 2.0: + tile ( size-expr-list ) */ + +static tree +c_parser_oacc_clause_tile (c_parser *parser, tree list) +{ + tree c, num = error_mark_node; + HOST_WIDE_INT n; + location_t loc; + tree tile = NULL_TREE; + vec *tvec = make_tree_vector (); + + check_no_duplicate_clause (list, OMP_CLAUSE_TILE, "tile"); + + loc = c_parser_peek_token (parser)->location; + if (!c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>")) + { + release_tree_vector (tvec); + return list; + } + + do + { + if (c_parser_next_token_is (parser, CPP_MULT)) + { + c_parser_consume_token (parser); + num = integer_minus_one_node; + } + else + { + num = c_parser_expr_no_commas (parser, NULL).value; + + if (num == error_mark_node) + { + c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, + "expected %<)%>"); + release_tree_vector (tvec); + return list; + } + + mark_exp_read (num); + num = c_fully_fold (num, false, NULL); + + if (!INTEGRAL_TYPE_P (TREE_TYPE (num)) + || !tree_fits_shwi_p (num) + || (n = tree_to_shwi (num)) <= 0 + || (int) n != n) + { + error_at (loc, + "tile argument needs positive constant integer " + "expression"); + release_tree_vector (tvec); + c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, + "expected %<)%>"); + return list; + } + } + + if (num == error_mark_node) + { + error_at (loc, "expected positive integer or %<)%>"); + release_tree_vector (tvec); + return list; + } + + vec_safe_push (tvec, num); + if (c_parser_next_token_is (parser, CPP_COMMA)) + c_parser_consume_token (parser); + } + while (c_parser_next_token_is_not (parser, CPP_CLOSE_PAREN)); + + /* Consume the trailing ')'. */ + c_parser_consume_token (parser); + + c = build_omp_clause (loc, OMP_CLAUSE_TILE); + tile = build_tree_list_vec (tvec); + OMP_CLAUSE_TILE_LIST (c) = tile; + OMP_CLAUSE_CHAIN (c) = list; + release_tree_vector (tvec); + return c; +} + +/* OpenACC 2.0: + use_device ( variable-list ) */ + +static tree +c_parser_oacc_clause_use_device (c_parser *parser, tree list) +{ + return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_USE_DEVICE, list); +} + /* OpenACC: wait ( int-expr-list ) */ @@ -10980,74 +11325,20 @@ c_parser_omp_clause_shared (c_parser *parser, tree list) } /* OpenMP 3.0: - untied */ + untied (FIXME: should we allow duplicates?) -static tree -c_parser_omp_clause_untied (c_parser *parser ATTRIBUTE_UNUSED, tree list) -{ - tree c; - - /* FIXME: Should we allow duplicates? */ - check_no_duplicate_clause (list, OMP_CLAUSE_UNTIED, "untied"); - - c = build_omp_clause (c_parser_peek_token (parser)->location, - OMP_CLAUSE_UNTIED); - OMP_CLAUSE_CHAIN (c) = list; - - return c; -} - -/* OpenACC: - vector_length ( expression ) */ - -static tree -c_parser_omp_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, - "% 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: + OpenMP 4.0: inbranch - notinbranch */ + notinbranch + + OpenACC 2.0: + auto + independent + nohost + seq */ static tree -c_parser_omp_clause_branch (c_parser *parser ATTRIBUTE_UNUSED, +c_parser_omp_simple_clause (c_parser *parser ATTRIBUTE_UNUSED, enum omp_clause_code code, tree list) { check_no_duplicate_clause (list, code, omp_clause_code_name[code]); @@ -11579,14 +11870,17 @@ c_parser_omp_clause_uniform (c_parser *parser, tree list) } /* Parse all OpenACC clauses. The set clauses allowed by the directive - is a bitmask in MASK. Return the list of clauses found. */ + is a bitmask in MASK. DTYPE_MASK denotes which clauses may follow a + device_type clause. Return the list of clauses found. */ -static tree +tree c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, - const char *where, bool finish_p = true) + const char *where, omp_clause_mask dtype_mask = 0, + bool finish_p = true, bool scan_dtype = true) { tree clauses = NULL; bool first = true; + bool seen_dtype = false; while (c_parser_next_token_is_not (parser, CPP_PRAGMA_EOL)) { @@ -11598,15 +11892,35 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, if (!first && c_parser_next_token_is (parser, CPP_COMMA)) c_parser_consume_token (parser); + if (!scan_dtype && c_parser_omp_clause_name (parser, false) + == PRAGMA_OACC_CLAUSE_DEVICE_TYPE) + return clauses; + here = c_parser_peek_token (parser)->location; c_kind = c_parser_omp_clause_name (parser); + if (seen_dtype && c_kind != PRAGMA_OMP_CLAUSE_NONE + && c_kind != PRAGMA_OACC_CLAUSE_DEVICE_TYPE) + { + error_at (here, "invalid clauses following device_type"); + goto saw_error; + } + switch (c_kind) { case PRAGMA_OACC_CLAUSE_ASYNC: clauses = c_parser_oacc_clause_async (parser, clauses); c_name = "async"; break; + case PRAGMA_OACC_CLAUSE_AUTO: + clauses = c_parser_omp_simple_clause (parser, OMP_CLAUSE_AUTO, + clauses); + c_name = "auto"; + break; + case PRAGMA_OACC_CLAUSE_BIND: + clauses = c_parser_oacc_clause_bind (parser, clauses); + c_name = "bind"; + break; case PRAGMA_OACC_CLAUSE_COLLAPSE: clauses = c_parser_omp_clause_collapse (parser, clauses); c_name = "collapse"; @@ -11631,10 +11945,24 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); c_name = "delete"; break; + case PRAGMA_OMP_CLAUSE_DEFAULT: + clauses = c_parser_omp_clause_default (parser, clauses, true); + c_name = "default"; + break; case PRAGMA_OACC_CLAUSE_DEVICE: clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); c_name = "device"; break; + case PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "device_resident"; + break; + case PRAGMA_OACC_CLAUSE_DEVICE_TYPE: + clauses = c_parser_oacc_clause_device_type (parser, dtype_mask, + clauses); + c_name = "device_type"; + seen_dtype = true; + break; case PRAGMA_OACC_CLAUSE_DEVICEPTR: clauses = c_parser_oacc_data_clause_deviceptr (parser, clauses); c_name = "deviceptr"; @@ -11643,6 +11971,11 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, clauses = c_parser_omp_clause_firstprivate (parser, clauses); c_name = "firstprivate"; break; + case PRAGMA_OACC_CLAUSE_GANG: + c_name = "gang"; + clauses = c_parser_oacc_shape_clause (parser, c_kind, c_name, + clauses); + break; case PRAGMA_OACC_CLAUSE_HOST: clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); c_name = "host"; @@ -11651,13 +11984,29 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, clauses = c_parser_omp_clause_if (parser, clauses); c_name = "if"; break; + case PRAGMA_OACC_CLAUSE_INDEPENDENT: + clauses = c_parser_omp_simple_clause (parser, OMP_CLAUSE_INDEPENDENT, + clauses); + c_name = "independent"; + break; + case PRAGMA_OACC_CLAUSE_LINK: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "link"; + break; + case PRAGMA_OACC_CLAUSE_NOHOST: + clauses = c_parser_omp_simple_clause (parser, OMP_CLAUSE_NOHOST, + clauses); + c_name = "nohost"; + break; case PRAGMA_OACC_CLAUSE_NUM_GANGS: - clauses = c_parser_omp_clause_num_gangs (parser, clauses); c_name = "num_gangs"; + clauses = c_parser_omp_positive_int_clause (parser, c_kind, c_name, + clauses); break; case PRAGMA_OACC_CLAUSE_NUM_WORKERS: - clauses = c_parser_omp_clause_num_workers (parser, clauses); c_name = "num_workers"; + clauses = c_parser_omp_positive_int_clause (parser, c_kind, c_name, + clauses); break; case PRAGMA_OACC_CLAUSE_PRESENT: clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); @@ -11687,18 +12036,38 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, 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"; + case PRAGMA_OACC_CLAUSE_SEQ: + clauses = c_parser_omp_simple_clause (parser, OMP_CLAUSE_SEQ, + clauses); + c_name = "seq"; + break; + case PRAGMA_OACC_CLAUSE_TILE: + clauses = c_parser_oacc_clause_tile (parser, clauses); + c_name = "tile"; + break; + case PRAGMA_OACC_CLAUSE_USE_DEVICE: + clauses = c_parser_oacc_clause_use_device (parser, clauses); + c_name = "use_device"; + break; + case PRAGMA_OACC_CLAUSE_VECTOR: + c_name = "vector"; + clauses = c_parser_oacc_shape_clause (parser, c_kind, c_name, + clauses); break; case PRAGMA_OACC_CLAUSE_VECTOR_LENGTH: - clauses = c_parser_omp_clause_vector_length (parser, clauses); c_name = "vector_length"; + clauses = c_parser_omp_positive_int_clause (parser, c_kind, c_name, + clauses); break; case PRAGMA_OACC_CLAUSE_WAIT: clauses = c_parser_oacc_clause_wait (parser, clauses); c_name = "wait"; break; + case PRAGMA_OACC_CLAUSE_WORKER: + c_name = "worker"; + clauses = c_parser_oacc_shape_clause (parser, c_kind, c_name, + clauses); + break; default: c_parser_error (parser, "expected %<#pragma acc%> clause"); goto saw_error; @@ -11715,11 +12084,17 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, } } + if (!scan_dtype) + return clauses; + saw_error: c_parser_skip_to_pragma_eol (parser); if (finish_p) - return c_finish_omp_clauses (clauses); + { + clauses = oacc_filter_device_types (clauses); + return c_finish_omp_clauses (clauses, true); + } return clauses; } @@ -11790,8 +12165,9 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask, c_name = "nowait"; break; case PRAGMA_OMP_CLAUSE_NUM_THREADS: - clauses = c_parser_omp_clause_num_threads (parser, clauses); c_name = "num_threads"; + clauses = c_parser_omp_positive_int_clause (parser, c_kind, c_name, + clauses); break; case PRAGMA_OMP_CLAUSE_ORDERED: clauses = c_parser_omp_clause_ordered (parser, clauses); @@ -11814,18 +12190,19 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask, c_name = "shared"; break; case PRAGMA_OMP_CLAUSE_UNTIED: - clauses = c_parser_omp_clause_untied (parser, clauses); + clauses = c_parser_omp_simple_clause (parser, OMP_CLAUSE_UNTIED, + clauses); c_name = "untied"; break; case PRAGMA_OMP_CLAUSE_INBRANCH: case PRAGMA_CILK_CLAUSE_MASK: - clauses = c_parser_omp_clause_branch (parser, OMP_CLAUSE_INBRANCH, + clauses = c_parser_omp_simple_clause (parser, OMP_CLAUSE_INBRANCH, clauses); c_name = "inbranch"; break; case PRAGMA_OMP_CLAUSE_NOTINBRANCH: case PRAGMA_CILK_CLAUSE_NOMASK: - clauses = c_parser_omp_clause_branch (parser, OMP_CLAUSE_NOTINBRANCH, + clauses = c_parser_omp_simple_clause (parser, OMP_CLAUSE_NOTINBRANCH, clauses); c_name = "notinbranch"; break; @@ -11948,7 +12325,7 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask, c_parser_skip_to_pragma_eol (parser); if (finish_p) - return c_finish_omp_clauses (clauses); + return c_finish_omp_clauses (clauses, false); return clauses; } @@ -11971,8 +12348,6 @@ c_parser_omp_structured_block (c_parser *parser) /* OpenACC 2.0: # pragma acc cache (variable-list) new-line - - LOC is the location of the #pragma token. */ static tree @@ -11981,7 +12356,7 @@ c_parser_oacc_cache (location_t loc, c_parser *parser) tree stmt, clauses; clauses = c_parser_omp_var_list_parens (parser, OMP_CLAUSE__CACHE_, NULL); - clauses = c_finish_omp_clauses (clauses); + clauses = c_finish_omp_clauses (clauses, true); c_parser_skip_to_pragma_eol (parser); @@ -11997,8 +12372,6 @@ c_parser_oacc_cache (location_t loc, c_parser *parser) /* 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 \ @@ -12020,7 +12393,8 @@ 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"); + "#pragma acc data", + OACC_DATA_CLAUSE_MASK); block = c_begin_omp_parallel (); add_stmt (c_parser_omp_structured_block (parser)); @@ -12031,57 +12405,190 @@ c_parser_oacc_data (location_t loc, c_parser *parser) } /* OpenACC 2.0: - # pragma acc kernels oacc-kernels-clause[optseq] new-line - structured-block - - LOC is the location of the #pragma token. + # pragma acc declare oacc-data-clause[optseq] new-line */ -#define OACC_KERNELS_CLAUSE_MASK \ - ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \ +#define OACC_DECLARE_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_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_OACC_CLAUSE_IF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_LINK) \ | (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) ) + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) ) + +static void +c_parser_oacc_declare (c_parser *parser) +{ + location_t pragma_loc = c_parser_peek_token (parser)->location; + tree clauses; + + c_parser_consume_pragma (parser); + + clauses = c_parser_oacc_all_clauses (parser, OACC_DECLARE_CLAUSE_MASK, + "#pragma acc declare"); + if (!clauses) + { + error_at (pragma_loc, + "no valid clauses specified in %<#pragma acc declare%>"); + return; + } + for (tree t = clauses; t; t = OMP_CLAUSE_CHAIN (t)) + { + location_t loc = OMP_CLAUSE_LOCATION (t); + tree decl = OMP_CLAUSE_DECL (t); + if (!DECL_P (decl)) + { + error_at (loc, "subarray in %<#pragma acc declare%>"); + continue; + } + gcc_assert (OMP_CLAUSE_CODE (t) == OMP_CLAUSE_MAP); + switch (OMP_CLAUSE_MAP_KIND (t)) + { + case GOMP_MAP_FORCE_ALLOC: + case GOMP_MAP_FORCE_TO: + case GOMP_MAP_FORCE_DEVICEPTR: + case GOMP_MAP_DEVICE_RESIDENT: + break; + + case GOMP_MAP_POINTER: + /* Generated by c_finish_omp_clauses from array sections; + avoid spurious diagnostics. */ + break; + + case GOMP_MAP_LINK: + if (!global_bindings_p () && !DECL_EXTERNAL (decl)) + { + error_at (loc, + "invalid variable %qD in %<#pragma acc declare link%>", + decl); + continue; + } + break; + + default: + if (global_bindings_p ()) + { + error_at (loc, "invalid OpenACC clause at file scope"); + continue; + } + if (DECL_EXTERNAL (decl)) + { + error_at (loc, + "invalid use of % variable %qD " + "in %<#pragma acc declare%>", decl); + continue; + } + break; + } + + /* Store the clause in an attribute on the variable, at file + scope, or the function, at block scope. */ + tree decl_for_attr; + if (global_bindings_p ()) + { + decl_for_attr = decl; + tree prev_attr = lookup_attribute ("oacc declare", + DECL_ATTRIBUTES (decl)); + if (prev_attr) + { + tree p = TREE_VALUE (prev_attr); + error_at (loc, + "variable %qD used more than once with " + "%<#pragma acc declare%>", decl); + inform (OMP_CLAUSE_LOCATION (TREE_VALUE (p)), + "previous directive was here"); + continue; + } + } + else + { + bool ok = true; + decl_for_attr = current_function_decl; + tree prev_attr = lookup_attribute ("oacc declare", + DECL_ATTRIBUTES (decl_for_attr)); + for (; + prev_attr; + prev_attr = lookup_attribute ("oacc declare", + TREE_CHAIN (prev_attr))) + { + tree p = TREE_VALUE (prev_attr); + tree cl = TREE_VALUE (p); + if (OMP_CLAUSE_DECL (cl) == decl) + { + error_at (loc, + "variable %qD used more than once with " + "%<#pragma acc declare%>", decl); + inform (OMP_CLAUSE_LOCATION (cl), + "previous directive was here"); + ok = false; + break; + } + } + if (!ok) + continue; + } + tree attr = tree_cons (NULL_TREE, t, NULL_TREE); + tree attrs = tree_cons (get_identifier ("oacc declare"), + attr, NULL_TREE); + decl_attributes (&decl_for_attr, attrs, 0); + } +} + +/* Split the 'clauses' into a set of 'loop' clauses and a set of + 'not-loop' clauses. */ static tree -c_parser_oacc_kernels (location_t loc, c_parser *parser, char *p_name) +oacc_split_loop_clauses (tree clauses, tree *not_loop_clauses) { - tree stmt, clauses = NULL_TREE, block; + tree loop_clauses, next, c; - strcat (p_name, " kernels"); + loop_clauses = *not_loop_clauses = NULL_TREE; - if (c_parser_next_token_is (parser, CPP_NAME)) + for (; clauses ; clauses = next) { - 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; + next = OMP_CLAUSE_CHAIN (clauses); + + switch (OMP_CLAUSE_CODE (clauses)) + { + case OMP_CLAUSE_COLLAPSE: + case OMP_CLAUSE_REDUCTION: + case OMP_CLAUSE_GANG: + case OMP_CLAUSE_VECTOR: + case OMP_CLAUSE_WORKER: + case OMP_CLAUSE_AUTO: + case OMP_CLAUSE_SEQ: + OMP_CLAUSE_CHAIN (clauses) = loop_clauses; + loop_clauses = clauses; + break; + + case OMP_CLAUSE_PRIVATE: + c = build_omp_clause (OMP_CLAUSE_LOCATION (clauses), + OMP_CLAUSE_CODE (clauses)); + OMP_CLAUSE_DECL (c) = OMP_CLAUSE_DECL (clauses); + OMP_CLAUSE_CHAIN (c) = loop_clauses; + loop_clauses = c; + /* FALL THROUGH */ + + default: + OMP_CLAUSE_CHAIN (clauses) = *not_loop_clauses; + *not_loop_clauses = clauses; + break; } } - clauses = c_parser_oacc_all_clauses (parser, OACC_KERNELS_CLAUSE_MASK, - p_name); + if (*not_loop_clauses) + c_finish_omp_clauses (*not_loop_clauses, true); - block = c_begin_omp_parallel (); - add_stmt (c_parser_omp_structured_block (parser)); + if (loop_clauses) + c_finish_omp_clauses (loop_clauses, true); - stmt = c_finish_oacc_kernels (loc, clauses, block); - - return stmt; + return loop_clauses; } /* OpenACC 2.0: @@ -12090,9 +12597,6 @@ c_parser_oacc_kernels (location_t loc, c_parser *parser, char *p_name) 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 \ @@ -12116,28 +12620,26 @@ 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)) + 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; + p = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value); + c_parser_consume_token (parser); } - const char *p = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value); if (strcmp (p, "data") != 0) { - c_parser_error (parser, "invalid pragma"); + error_at (loc, enter + ? "expected % after %<#pragma acc enter%>" + : "expected % after %<#pragma acc exit%>"); + parser->error = true; 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"); @@ -12160,27 +12662,72 @@ c_parser_oacc_enter_exit_data (c_parser *parser, bool enter) add_stmt (stmt); } +/* OpenACC 2.0: + # pragma acc host_data oacc-data-clause[optseq] new-line + structured-block +*/ + +#define OACC_HOST_DATA_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_USE_DEVICE) ) + +static tree +c_parser_oacc_host_data (location_t loc, c_parser *parser) +{ + tree stmt, clauses, block; + + clauses = c_parser_oacc_all_clauses (parser, OACC_HOST_DATA_CLAUSE_MASK, + "#pragma acc host_data"); + + block = c_begin_omp_parallel (); + add_stmt (c_parser_omp_structured_block (parser)); + stmt = c_finish_oacc_host_data (loc, clauses, block); + return 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_OACC_CLAUSE_COLLAPSE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_TYPE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_GANG) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WORKER) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_AUTO) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_INDEPENDENT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_TILE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) ) +#define OACC_LOOP_CLAUSE_DEVICE_TYPE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COLLAPSE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_GANG) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WORKER) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_AUTO) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_INDEPENDENT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_TILE) ) + static tree -c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name) +c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name, + omp_clause_mask mask, tree *cclauses) { tree stmt, clauses, block; strcat (p_name, " loop"); + mask |= OACC_LOOP_CLAUSE_MASK; - clauses = c_parser_oacc_all_clauses (parser, OACC_LOOP_CLAUSE_MASK, p_name); + clauses = c_parser_oacc_all_clauses (parser, mask, p_name, + OACC_LOOP_CLAUSE_DEVICE_TYPE_MASK, + cclauses == NULL); + if (cclauses) + clauses = oacc_split_loop_clauses (clauses, cclauses); block = c_begin_compound_stmt (true); stmt = c_parser_omp_for_loop (loc, parser, OACC_LOOP, clauses, NULL); @@ -12191,10 +12738,68 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name) } /* 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_OACC_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_DEFAULT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_TYPE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_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) ) + +#define OACC_KERNELS_CLAUSE_DEVICE_TYPE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \ + | (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, block; + omp_clause_mask mask; + + strcat (p_name, " kernels"); + + mask = OACC_KERNELS_CLAUSE_MASK; + if (c_parser_next_token_is (parser, CPP_NAME)) + { + stmt = c_parser_peek_token (parser)->value; + if (!strcmp ("loop", IDENTIFIER_POINTER (stmt))) + { + tree kernel_clauses; + + c_parser_consume_token (parser); + mask |= OACC_LOOP_CLAUSE_MASK; + block = c_begin_omp_parallel (); + c_parser_oacc_loop (loc, parser, p_name, mask, &kernel_clauses); + stmt = c_finish_oacc_kernels (loc, kernel_clauses, block); + return stmt; + } + } + + clauses = c_parser_oacc_all_clauses (parser, mask, p_name, + OACC_KERNELS_CLAUSE_DEVICE_TYPE_MASK); + + 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 parallel oacc-parallel-clause[optseq] new-line structured-block - - LOC is the location of the #pragma token. */ #define OACC_PARALLEL_CLAUSE_MASK \ @@ -12203,8 +12808,11 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name) | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_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_DEFAULT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_TYPE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \ | (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) \ @@ -12216,48 +12824,227 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name) | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) +#define OACC_PARALLEL_CLAUSE_DEVICE_TYPE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \ + | (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_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; + tree stmt, clauses, block; + omp_clause_mask mask, dmask; strcat (p_name, " parallel"); + mask = OACC_PARALLEL_CLAUSE_MASK; + dmask = OACC_PARALLEL_CLAUSE_DEVICE_TYPE_MASK; 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) + stmt = c_parser_peek_token (parser)->value; + if (!strcmp ("loop", IDENTIFIER_POINTER (stmt))) { + tree parallel_clauses; + c_parser_consume_token (parser); + mask |= OACC_LOOP_CLAUSE_MASK; 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; + c_parser_oacc_loop (loc, parser, p_name, mask, ¶llel_clauses); + stmt = c_finish_oacc_parallel (loc, parallel_clauses, block); return stmt; } } - clauses = c_parser_oacc_all_clauses (parser, OACC_PARALLEL_CLAUSE_MASK, - p_name); + clauses = c_parser_oacc_all_clauses (parser, mask, p_name, dmask); 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 routine oacc-routine-clause[optseq] new-line + function-definition + + # pragma acc routine ( name ) oacc-routine-clause[optseq] new-line +*/ + +#define OACC_ROUTINE_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_TYPE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_GANG) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WORKER) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_BIND) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NOHOST)) + +#define OACC_ROUTINE_CLAUSE_DEVICE_TYPE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_GANG) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WORKER) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_BIND)) + +static void +c_parser_oacc_routine (c_parser *parser, enum pragma_context context) +{ + tree name = NULL_TREE; + location_t here = c_parser_peek_token (parser)->location; + + c_parser_consume_pragma (parser); + + /* Scan for optional '( name )'. */ + if (c_parser_peek_token (parser)->type == CPP_OPEN_PAREN) + { + c_parser_consume_token (parser); + + if (c_parser_next_token_is_not (parser, CPP_NAME) + || c_parser_peek_token (parser)->id_kind != C_ID_ID) + c_parser_error (parser, "expected identifier"); + + // name should be an IDENTIFIER_NODE + name = c_parser_peek_token (parser)->value; + + if (name == NULL_TREE) + { + undeclared_variable (c_parser_peek_token (parser)->location, + c_parser_peek_token (parser)->value); + name = error_mark_node; + } + + c_parser_consume_token (parser); + + if (name == error_mark_node) + return; + + c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, 0); + } + + /* Build a chain of clauses. */ + parser->in_pragma = true; + tree clauses = NULL_TREE; + clauses = c_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK, + "#pragma acc routine", + OACC_ROUTINE_CLAUSE_DEVICE_TYPE_MASK); + + /* Check of the presence if gang, worker, vector and seq clauses, and + throw an error if more than one of those clauses is specified. */ + int parallelism = 0; + tree c; + + for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) + switch (OMP_CLAUSE_CODE (c)) + { + case OMP_CLAUSE_GANG: + case OMP_CLAUSE_WORKER: + case OMP_CLAUSE_VECTOR: + case OMP_CLAUSE_SEQ: + ++parallelism; + break; + default: + break; + } + + if (parallelism > 1) + { + error_at (here, "invalid combination of gang, worker, vector or seq for" + "%<#pragma acc routine%>"); + } + + if (name) + { + TREE_CHAIN (name) = clauses; + vec_safe_push (parser->oacc_routines, name); + } + else + { + if (context != pragma_external) + { + c_parser_error (parser, "%<#pragma acc routine%> must be " + "followed by function declaration or definition"); + return; + } + + if (c_parser_next_token_is (parser, CPP_KEYWORD) + && c_parser_peek_token (parser)->keyword == RID_EXTENSION) + { + int ext = disable_extension_diagnostics (); + do + c_parser_consume_token (parser); + while (c_parser_next_token_is (parser, CPP_KEYWORD) + && c_parser_peek_token (parser)->keyword + == RID_EXTENSION); + c_parser_declaration_or_fndef (parser, true, true, true, false, + true, NULL, vNULL, clauses, true); + restore_extension_diagnostics (ext); + } + else + c_parser_declaration_or_fndef (parser, true, true, true, false, + true, NULL, vNULL, clauses, true); + } +} + +static void +c_finish_oacc_routine (c_parser *parser, tree fndecl, tree clauses, + bool named) +{ + if (fndecl == NULL_TREE || TREE_CODE (fndecl) != FUNCTION_DECL) + { + if (!named) + return; + + error ("%<#pragma acc routine%> not immediately followed by " + "a function declaration or definition"); + gcc_unreachable(); + return; + } + + if (!named) + { + bool found = false; + int i; + tree t; + + for (i = 0; vec_safe_iterate (parser->oacc_routines, i, &t); i++) + { + if (!strcmp (IDENTIFIER_POINTER (DECL_NAME (fndecl)), + IDENTIFIER_POINTER (t))) + { + found = true; + clauses = TREE_CHAIN (t); + break; + } + } + + if (!found) + return; + } + + if (clauses != NULL_TREE) + clauses = tree_cons (NULL_TREE, clauses, NULL_TREE); + clauses = build_tree_list (get_identifier ("omp declare target"), + clauses); + TREE_CHAIN (clauses) = DECL_ATTRIBUTES (fndecl); + DECL_ATTRIBUTES (fndecl) = clauses; +} + +/* 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_OACC_CLAUSE_DEVICE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_TYPE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_HOST) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SELF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) + +#define OACC_UPDATE_CLAUSE_DEVICE_TYPE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) static void @@ -12268,7 +13055,8 @@ c_parser_oacc_update (c_parser *parser) c_parser_consume_pragma (parser); tree clauses = c_parser_oacc_all_clauses (parser, OACC_UPDATE_CLAUSE_MASK, - "#pragma acc update"); + "#pragma acc update", + OACC_UPDATE_CLAUSE_DEVICE_TYPE_MASK); if (find_omp_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE) { error_at (loc, @@ -12289,8 +13077,6 @@ c_parser_oacc_update (c_parser *parser) /* 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 \ @@ -12844,7 +13630,7 @@ c_parser_omp_for_loop (location_t loc, c_parser *parser, enum tree_code code, if (i > 0) vec_safe_push (for_block, c_begin_compound_stmt (true)); c_parser_declaration_or_fndef (parser, true, true, true, true, true, - NULL, vNULL); + NULL, vNULL, NULL_TREE, false); decl = check_for_loop_decls (for_loc, flag_isoc99); if (decl == NULL) goto error_init; @@ -13115,7 +13901,7 @@ omp_split_clauses (location_t loc, enum tree_code code, c_omp_split_clauses (loc, code, mask, clauses, cclauses); for (i = 0; i < C_OMP_CLAUSE_SPLIT_COUNT; i++) if (cclauses[i]) - cclauses[i] = c_finish_omp_clauses (cclauses[i]); + cclauses[i] = c_finish_omp_clauses (cclauses[i], false); } /* OpenMP 4.0: @@ -14032,12 +14818,12 @@ c_parser_omp_declare_simd (c_parser *parser, enum pragma_context context) while (c_parser_next_token_is (parser, CPP_KEYWORD) && c_parser_peek_token (parser)->keyword == RID_EXTENSION); c_parser_declaration_or_fndef (parser, true, true, true, false, true, - NULL, clauses); + NULL, clauses, NULL_TREE, false); restore_extension_diagnostics (ext); } else c_parser_declaration_or_fndef (parser, true, true, true, false, true, - NULL, clauses); + NULL, clauses, NULL_TREE, false); break; case pragma_struct: case pragma_param: @@ -14057,7 +14843,8 @@ c_parser_omp_declare_simd (c_parser *parser, enum pragma_context context) if (c_parser_next_tokens_start_declaration (parser)) { c_parser_declaration_or_fndef (parser, true, true, true, true, - true, NULL, clauses); + true, NULL, clauses, NULL_TREE, + false); restore_extension_diagnostics (ext); break; } @@ -14066,7 +14853,7 @@ c_parser_omp_declare_simd (c_parser *parser, enum pragma_context context) else if (c_parser_next_tokens_start_declaration (parser)) { c_parser_declaration_or_fndef (parser, true, true, true, true, true, - NULL, clauses); + NULL, clauses, NULL_TREE, false); break; } c_parser_error (parser, "%<#pragma omp declare simd%> must be followed by " @@ -14634,6 +15421,9 @@ c_parser_omp_construct (c_parser *parser) switch (p_kind) { + case PRAGMA_OACC_ATOMIC: + c_parser_omp_atomic (loc, parser); + return; case PRAGMA_OACC_CACHE: strcpy (p_name, "#pragma acc"); stmt = c_parser_oacc_cache (loc, parser); @@ -14641,13 +15431,16 @@ c_parser_omp_construct (c_parser *parser) case PRAGMA_OACC_DATA: stmt = c_parser_oacc_data (loc, parser); break; + case PRAGMA_OACC_HOST_DATA: + stmt = c_parser_oacc_host_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); + stmt = c_parser_oacc_loop (loc, parser, p_name, mask, NULL); break; case PRAGMA_OACC_PARALLEL: strcpy (p_name, "#pragma acc"); @@ -15100,7 +15893,7 @@ c_parser_cilk_for (c_parser *parser, tree grain) tree clauses = build_omp_clause (EXPR_LOCATION (grain), OMP_CLAUSE_SCHEDULE); OMP_CLAUSE_SCHEDULE_KIND (clauses) = OMP_CLAUSE_SCHEDULE_CILKFOR; OMP_CLAUSE_SCHEDULE_CHUNK_EXPR (clauses) = grain; - clauses = c_finish_omp_clauses (clauses); + clauses = c_finish_omp_clauses (clauses, false); tree block = c_begin_compound_stmt (true); tree sb = push_stmt_list (); @@ -15165,7 +15958,7 @@ c_parser_cilk_for (c_parser *parser, tree grain) OMP_CLAUSE_OPERAND (c, 0) = cilk_for_number_of_iterations (omp_for); OMP_CLAUSE_CHAIN (c) = clauses; - OMP_PARALLEL_CLAUSES (omp_par) = c_finish_omp_clauses (c); + OMP_PARALLEL_CLAUSES (omp_par) = c_finish_omp_clauses (c, false); add_stmt (omp_par); } @@ -15408,6 +16201,8 @@ c_parse_file (void) if (tparser.tokens == &tparser.tokens_buf[0]) the_parser->tokens = &the_parser->tokens_buf[0]; + the_parser->oacc_routines = NULL; + /* Initialize EH, if we've been told to do so. */ if (flag_exceptions) using_eh_for_cleanups (); diff --git gcc/c/c-tree.h gcc/c/c-tree.h index 7a72665..8750bd7 100644 --- gcc/c/c-tree.h +++ gcc/c/c-tree.h @@ -643,13 +643,14 @@ 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_finish_oacc_host_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); extern tree c_finish_omp_task (location_t, tree, tree); extern void c_finish_omp_cancel (location_t, tree); extern void c_finish_omp_cancellation_point (location_t, tree); -extern tree c_finish_omp_clauses (tree); +extern tree c_finish_omp_clauses (tree, bool); extern tree c_build_va_arg (location_t, tree, tree); extern tree c_finish_transaction (location_t, tree, int); extern bool c_tree_equal (tree, tree); diff --git gcc/c/c-typeck.c gcc/c/c-typeck.c index 91735b5..e27a1c7 100644 --- gcc/c/c-typeck.c +++ gcc/c/c-typeck.c @@ -11449,6 +11449,25 @@ c_finish_oacc_data (location_t loc, tree clauses, tree block) return add_stmt (stmt); } +/* Generate OACC_HOST_DATA, with CLAUSES and BLOCK as its compound + statement. LOC is the location of the OACC_HOST_DATA. */ + +tree +c_finish_oacc_host_data (location_t loc, tree clauses, tree block) +{ + tree stmt; + + block = c_end_compound_stmt (loc, block, true); + + stmt = make_node (OACC_HOST_DATA); + TREE_TYPE (stmt) = void_type_node; + OACC_HOST_DATA_CLAUSES (stmt) = clauses; + OACC_HOST_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 @@ -12048,13 +12067,14 @@ c_find_omp_placeholder_r (tree *tp, int *, void *data) Remove any elements from the list that are invalid. */ tree -c_finish_omp_clauses (tree clauses) +c_finish_omp_clauses (tree clauses, bool oacc) { bitmap_head generic_head, firstprivate_head, lastprivate_head; - bitmap_head aligned_head; + bitmap_head aligned_head, oacc_data_head; tree c, t, *pc; bool branch_seen = false; bool copyprivate_seen = false; + bool oacc_data = false; tree *nowait_clause = NULL; bitmap_obstack_initialize (NULL); @@ -12062,6 +12082,7 @@ c_finish_omp_clauses (tree clauses) bitmap_initialize (&firstprivate_head, &bitmap_default_obstack); bitmap_initialize (&lastprivate_head, &bitmap_default_obstack); bitmap_initialize (&aligned_head, &bitmap_default_obstack); + bitmap_initialize (&oacc_data_head, &bitmap_default_obstack); for (pc = &clauses, c = clauses; c ; c = *pc) { @@ -12077,11 +12098,16 @@ c_finish_omp_clauses (tree clauses) case OMP_CLAUSE_PRIVATE: need_complete = true; + oacc_data = true; need_implicitly_determined = true; - goto check_dup_generic; + if (oacc) + goto check_dup_oacc; + else + goto check_dup_generic; case OMP_CLAUSE_REDUCTION: need_implicitly_determined = true; + oacc_data = false; t = OMP_CLAUSE_DECL (c); if (OMP_CLAUSE_REDUCTION_PLACEHOLDER (c) == NULL_TREE && (FLOAT_TYPE_P (TREE_TYPE (t)) @@ -12201,7 +12227,10 @@ c_finish_omp_clauses (tree clauses) OMP_CLAUSE_REDUCTION_INIT (c), NULL_TREE); TREE_SIDE_EFFECTS (OMP_CLAUSE_REDUCTION_INIT (c)) = 1; } - goto check_dup_generic; + if (oacc) + goto check_dup_oacc; + else + goto check_dup_generic; case OMP_CLAUSE_COPYPRIVATE: copyprivate_seen = true; @@ -12262,9 +12291,9 @@ c_finish_omp_clauses (tree clauses) omp_clause_code_name[OMP_CLAUSE_CODE (c)]); remove = true; } - else if (bitmap_bit_p (&generic_head, DECL_UID (t)) - || bitmap_bit_p (&firstprivate_head, DECL_UID (t)) - || bitmap_bit_p (&lastprivate_head, DECL_UID (t))) + if (bitmap_bit_p (&generic_head, DECL_UID (t)) + || bitmap_bit_p (&firstprivate_head, DECL_UID (t)) + || bitmap_bit_p (&lastprivate_head, DECL_UID (t))) { error_at (OMP_CLAUSE_LOCATION (c), "%qE appears more than once in data clauses", t); @@ -12274,6 +12303,39 @@ c_finish_omp_clauses (tree clauses) bitmap_set_bit (&generic_head, DECL_UID (t)); break; + check_dup_oacc: + t = OMP_CLAUSE_DECL (c); + if (TREE_CODE (t) != VAR_DECL && TREE_CODE (t) != PARM_DECL) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qE is not a variable in clause %qs", t, + omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + remove = true; + } + if (oacc_data) + { + if (bitmap_bit_p (&oacc_data_head, DECL_UID (t))) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qE appears more than once in data clauses", t); + remove = true; + } + else + bitmap_set_bit (&oacc_data_head, DECL_UID (t)); + } + else + { + if (bitmap_bit_p (&generic_head, DECL_UID (t))) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qE appears more than once in data clauses", t); + remove = true; + } + else + bitmap_set_bit (&generic_head, DECL_UID (t)); + } + break; + case OMP_CLAUSE_FIRSTPRIVATE: t = OMP_CLAUSE_DECL (c); need_complete = true; @@ -12284,15 +12346,29 @@ c_finish_omp_clauses (tree clauses) "%qE is not a variable in clause %", t); remove = true; } - else if (bitmap_bit_p (&generic_head, DECL_UID (t)) - || bitmap_bit_p (&firstprivate_head, DECL_UID (t))) + else if (oacc) { - error_at (OMP_CLAUSE_LOCATION (c), - "%qE appears more than once in data clauses", t); - remove = true; + if (bitmap_bit_p (&oacc_data_head, DECL_UID (t))) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qE appears more than once in data clauses", t); + remove = true; + } + else + bitmap_set_bit (&oacc_data_head, DECL_UID (t)); } else - bitmap_set_bit (&firstprivate_head, DECL_UID (t)); + { + if (bitmap_bit_p (&generic_head, DECL_UID (t)) + || bitmap_bit_p (&firstprivate_head, DECL_UID (t))) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qE appears more than once in data clauses", t); + remove = true; + } + else + bitmap_set_bit (&firstprivate_head, DECL_UID (t)); + } break; case OMP_CLAUSE_LASTPRIVATE: @@ -12415,7 +12491,8 @@ c_finish_omp_clauses (tree clauses) omp_clause_code_name[OMP_CLAUSE_CODE (c)]); remove = true; } - else if (bitmap_bit_p (&generic_head, DECL_UID (t))) + if ((oacc && bitmap_bit_p (&oacc_data_head, DECL_UID (t))) + || bitmap_bit_p (&generic_head, DECL_UID (t))) { if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) error ("%qD appears more than once in motion clauses", t); @@ -12423,6 +12500,8 @@ c_finish_omp_clauses (tree clauses) error ("%qD appears more than once in map clauses", t); remove = true; } + else if (oacc) + bitmap_set_bit (&oacc_data_head, DECL_UID (t)); else bitmap_set_bit (&generic_head, DECL_UID (t)); break; @@ -12482,10 +12561,15 @@ c_finish_omp_clauses (tree clauses) case OMP_CLAUSE_ASYNC: case OMP_CLAUSE_WAIT: case OMP_CLAUSE_AUTO: + case OMP_CLAUSE_INDEPENDENT: case OMP_CLAUSE_SEQ: case OMP_CLAUSE_GANG: case OMP_CLAUSE_WORKER: case OMP_CLAUSE_VECTOR: + case OMP_CLAUSE_USE_DEVICE: + case OMP_CLAUSE_BIND: + case OMP_CLAUSE_NOHOST: + case OMP_CLAUSE_TILE: pc = &OMP_CLAUSE_CHAIN (c); continue; diff --git gcc/cp/cp-gimplify.c gcc/cp/cp-gimplify.c index 70645b5..569733c 100644 --- gcc/cp/cp-gimplify.c +++ gcc/cp/cp-gimplify.c @@ -1533,7 +1533,8 @@ cxx_omp_clause_default_ctor (tree clause, tree decl, tree /*outer*/) tree cxx_omp_clause_copy_ctor (tree clause, tree dst, tree src) { - tree info = CP_OMP_CLAUSE_INFO (clause); + tree info = OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP ? NULL + : CP_OMP_CLAUSE_INFO (clause); tree ret = NULL; if (info) diff --git gcc/cp/cp-tree.h gcc/cp/cp-tree.h index 2a904a5..251ed38 100644 --- gcc/cp/cp-tree.h +++ gcc/cp/cp-tree.h @@ -5986,11 +5986,12 @@ extern void note_decl_for_pch (tree); extern tree omp_reduction_id (enum tree_code, tree, tree); extern tree cp_remove_omp_priv_cleanup_stmt (tree *, int *, void *); extern void cp_check_omp_declare_reduction (tree); -extern tree finish_omp_clauses (tree); +extern tree finish_omp_clauses (tree, bool); extern void finish_omp_threadprivate (tree); extern tree begin_omp_structured_block (void); extern tree finish_omp_structured_block (tree); extern tree finish_oacc_data (tree, tree); +extern tree finish_oacc_host_data (tree, tree); extern tree finish_oacc_kernels (tree, tree); extern tree finish_oacc_parallel (tree, tree); extern tree begin_omp_parallel (void); diff --git gcc/cp/parser.c gcc/cp/parser.c index cfb512b..6e177f6 100644 --- gcc/cp/parser.c +++ gcc/cp/parser.c @@ -1303,7 +1303,8 @@ cp_token_cache_new (cp_token *first, cp_token *last) } /* Diagnose if #pragma omp declare simd isn't followed immediately - by function declaration or definition. */ + by function declaration or definition. Likewise for + #pragma acc routine. */ static inline void cp_ensure_no_omp_declare_simd (cp_parser *parser) @@ -1314,6 +1315,13 @@ cp_ensure_no_omp_declare_simd (cp_parser *parser) "function declaration or definition"); parser->omp_declare_simd = NULL; } + + if (parser->oacc_routine && !parser->oacc_routine->error_seen) + { + error ("%<#pragma acc routine%> not immediately followed by " + "function declaration or definition"); + parser->oacc_routine = NULL; + } } /* Finalize #pragma omp declare simd clauses after FNDECL has been parsed, @@ -1336,6 +1344,58 @@ cp_finalize_omp_declare_simd (cp_parser *parser, tree fndecl) } } } + +/* Finalize #pragma acc routine clauses after FNDECL has been parsed, + and put that into "acc routine" attribute. */ + +static inline void +cp_finalize_oacc_routine (cp_parser *parser, tree fndecl) +{ + if (__builtin_expect (parser->omp_declare_simd != NULL, 0)) + { + if (fndecl == error_mark_node) + { + parser->omp_declare_simd = NULL; + return; + } + if (TREE_CODE (fndecl) != FUNCTION_DECL) + { + cp_ensure_no_omp_declare_simd (parser); + return; + } + } + else // Is this fndecl associated with a named routine? + { + if (fndecl == NULL_TREE || fndecl == error_mark_node + || TREE_CODE (fndecl) != FUNCTION_DECL) + return; + + bool found = false; + int i; + tree t, clauses = NULL_TREE; + + for (i = 0; vec_safe_iterate (parser->named_oacc_routines, i, &t); i++) + { + if (!strcmp (IDENTIFIER_POINTER (DECL_NAME (fndecl)), + IDENTIFIER_POINTER (t))) + { + found = true; + clauses = TREE_CHAIN (t); + break; + } + } + + if (!found) + return; + + if (clauses != NULL_TREE) + clauses = tree_cons (NULL_TREE, clauses, NULL_TREE); + clauses = build_tree_list (get_identifier ("omp declare target"), + clauses); + TREE_CHAIN (clauses) = DECL_ATTRIBUTES (fndecl); + DECL_ATTRIBUTES (fndecl) = clauses; + } +} /* Decl-specifiers. */ @@ -2185,6 +2245,9 @@ static tree cp_parser_late_parsing_omp_declare_simd static tree cp_parser_late_parsing_cilk_simd_fn_info (cp_parser *, tree); +static tree cp_parser_late_parsing_oacc_routine + (cp_parser *, tree); + static tree synthesize_implicit_template_parm (cp_parser *); static tree finish_fully_implicit_template @@ -2541,6 +2604,11 @@ static bool cp_parser_array_designator_p static bool cp_parser_skip_to_closing_square_bracket (cp_parser *); +/* OpenACC routines. */ +static tree cp_parser_oacc_all_clauses (cp_parser *, omp_clause_mask, + const char *, cp_token *, + omp_clause_mask, bool, bool); + /* Returns nonzero if we are parsing tentatively. */ static inline bool @@ -3561,6 +3629,10 @@ cp_parser_new (void) parser->implicit_template_parms = 0; parser->implicit_template_scope = 0; + /* The list of OpenACC routines pragmas is unitialized. */ + parser->oacc_routine = NULL; + parser->named_oacc_routines = NULL; + return parser; } @@ -17150,6 +17222,7 @@ cp_parser_init_declarator (cp_parser* parser, range_for_decl_p? SD_INITIALIZED : is_initialized, attributes, prefix_attributes, &pushed_scope); cp_finalize_omp_declare_simd (parser, decl); + cp_finalize_oacc_routine (parser, decl); /* Adjust location of decl if declarator->id_loc is more appropriate: set, and decl wasn't merged with another decl, in which case its location would be different from input_location, and more accurate. */ @@ -17263,6 +17336,7 @@ cp_parser_init_declarator (cp_parser* parser, if (decl && TREE_CODE (decl) == FUNCTION_DECL) cp_parser_save_default_args (parser, decl); cp_finalize_omp_declare_simd (parser, decl); + cp_finalize_oacc_routine (parser, decl); } /* Finish processing the declaration. But, skip member @@ -18329,11 +18403,15 @@ cp_parser_late_return_type_opt (cp_parser* parser, cp_declarator *declarator, bool cilk_simd_fn_vector_p = (parser->cilk_simd_fn_info && declarator && declarator->kind == cdk_id); - + + bool oacc_routine_p = (parser->oacc_routine + && declarator && declarator->kind == cdk_id); + /* Peek at the next token. */ token = cp_lexer_peek_token (parser->lexer); /* A late-specified return type is indicated by an initial '->'. */ - if (token->type != CPP_DEREF && !(declare_simd_p || cilk_simd_fn_vector_p)) + if (token->type != CPP_DEREF && !(declare_simd_p || cilk_simd_fn_vector_p + || oacc_routine_p)) return NULL_TREE; tree save_ccp = current_class_ptr; @@ -18360,6 +18438,10 @@ cp_parser_late_return_type_opt (cp_parser* parser, cp_declarator *declarator, declarator->std_attributes = cp_parser_late_parsing_omp_declare_simd (parser, declarator->std_attributes); + if (oacc_routine_p) + declarator->std_attributes + = cp_parser_late_parsing_oacc_routine (parser, + declarator->std_attributes); if (quals >= 0) { @@ -21097,6 +21179,7 @@ cp_parser_member_declaration (cp_parser* parser) } cp_finalize_omp_declare_simd (parser, decl); + cp_finalize_oacc_routine (parser, decl); /* Reset PREFIX_ATTRIBUTES. */ while (attributes && TREE_CHAIN (attributes) != first_attribute) @@ -23349,6 +23432,9 @@ cp_parser_function_definition_from_specifiers_and_declarator { cp_finalize_omp_declare_simd (parser, current_function_decl); parser->omp_declare_simd = NULL; + + cp_finalize_oacc_routine (parser, current_function_decl); + parser->oacc_routine = NULL; } if (!success_p) @@ -23910,6 +23996,7 @@ cp_parser_save_member_function_body (cp_parser* parser, /* Create the FUNCTION_DECL. */ fn = grokmethod (decl_specifiers, declarator, attributes); cp_finalize_omp_declare_simd (parser, fn); + cp_finalize_oacc_routine (parser, fn); /* If something went badly wrong, bail out now. */ if (fn == error_mark_node) { @@ -27529,11 +27616,13 @@ cp_parser_objc_at_dynamic_declaration (cp_parser *parser) returned and the token is consumed. */ static pragma_omp_clause -cp_parser_omp_clause_name (cp_parser *parser) +cp_parser_omp_clause_name (cp_parser *parser, bool consume_token = true) { pragma_omp_clause result = PRAGMA_OMP_CLAUSE_NONE; - if (cp_lexer_next_token_is_keyword (parser->lexer, RID_IF)) + if (cp_lexer_next_token_is_keyword (parser->lexer, RID_AUTO)) + result = PRAGMA_OACC_CLAUSE_AUTO; + else if (cp_lexer_next_token_is_keyword (parser->lexer, RID_IF)) result = PRAGMA_OMP_CLAUSE_IF; else if (cp_lexer_next_token_is_keyword (parser->lexer, RID_DEFAULT)) result = PRAGMA_OMP_CLAUSE_DEFAULT; @@ -27556,6 +27645,10 @@ cp_parser_omp_clause_name (cp_parser *parser) else if (!strcmp ("async", p)) result = PRAGMA_OACC_CLAUSE_ASYNC; break; + case 'b': + if (!strcmp ("bind", p)) + result = PRAGMA_OACC_CLAUSE_BIND; + break; case 'c': if (!strcmp ("collapse", p)) result = PRAGMA_OMP_CLAUSE_COLLAPSE; @@ -27575,6 +27668,11 @@ cp_parser_omp_clause_name (cp_parser *parser) result = PRAGMA_OMP_CLAUSE_DEPEND; else if (!strcmp ("device", p)) result = PRAGMA_OMP_CLAUSE_DEVICE; + else if (!strcmp ("device_resident", p)) + result = PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT; + else if (!strcmp ("device_type", p) + || !strcmp ("dtype", p)) + result = PRAGMA_OACC_CLAUSE_DEVICE_TYPE; else if (!strcmp ("deviceptr", p)) result = PRAGMA_OACC_CLAUSE_DEVICEPTR; else if (!strcmp ("dist_schedule", p)) @@ -27592,15 +27690,23 @@ cp_parser_omp_clause_name (cp_parser *parser) if (!strcmp ("host", p)) result = PRAGMA_OACC_CLAUSE_HOST; break; + case 'g': + if (!strcmp ("gang", p)) + result = PRAGMA_OACC_CLAUSE_GANG; + break; case 'i': if (!strcmp ("inbranch", p)) result = PRAGMA_OMP_CLAUSE_INBRANCH; + else if (!strcmp ("independent", p)) + result = PRAGMA_OACC_CLAUSE_INDEPENDENT; break; case 'l': if (!strcmp ("lastprivate", p)) result = PRAGMA_OMP_CLAUSE_LASTPRIVATE; else if (!strcmp ("linear", p)) result = PRAGMA_OMP_CLAUSE_LINEAR; + else if (!strcmp ("link", p)) + result = PRAGMA_OACC_CLAUSE_LINK; break; case 'm': if (!strcmp ("map", p)) @@ -27615,6 +27721,8 @@ cp_parser_omp_clause_name (cp_parser *parser) result = PRAGMA_OMP_CLAUSE_NOTINBRANCH; else if (!strcmp ("nowait", p)) result = PRAGMA_OMP_CLAUSE_NOWAIT; + else if (!strcmp ("nohost", p)) + result = PRAGMA_OACC_CLAUSE_NOHOST; else if (flag_cilkplus && !strcmp ("nomask", p)) result = PRAGMA_CILK_CLAUSE_NOMASK; else if (!strcmp ("num_gangs", p)) @@ -27661,8 +27769,10 @@ cp_parser_omp_clause_name (cp_parser *parser) result = PRAGMA_OMP_CLAUSE_SCHEDULE; else if (!strcmp ("sections", p)) result = PRAGMA_OMP_CLAUSE_SECTIONS; - else if (!strcmp ("self", p)) - result = PRAGMA_OACC_CLAUSE_SELF; + else if (!strcmp ("self", p)) /* "self" is a synonym for "host". */ + result = PRAGMA_OACC_CLAUSE_HOST; + else if (!strcmp ("seq", p)) + result = PRAGMA_OACC_CLAUSE_SEQ; else if (!strcmp ("shared", p)) result = PRAGMA_OMP_CLAUSE_SHARED; else if (!strcmp ("simdlen", p)) @@ -27673,6 +27783,8 @@ cp_parser_omp_clause_name (cp_parser *parser) result = PRAGMA_OMP_CLAUSE_TASKGROUP; else if (!strcmp ("thread_limit", p)) result = PRAGMA_OMP_CLAUSE_THREAD_LIMIT; + else if (!strcmp ("tile", p)) + result = PRAGMA_OACC_CLAUSE_TILE; else if (!strcmp ("to", p)) result = PRAGMA_OMP_CLAUSE_TO; break; @@ -27681,9 +27793,13 @@ cp_parser_omp_clause_name (cp_parser *parser) result = PRAGMA_OMP_CLAUSE_UNIFORM; else if (!strcmp ("untied", p)) result = PRAGMA_OMP_CLAUSE_UNTIED; + else if (!strcmp ("use_device", p)) + result = PRAGMA_OACC_CLAUSE_USE_DEVICE; break; case 'v': - if (!strcmp ("vector_length", p)) + if (!strcmp ("vector", p)) + result = PRAGMA_OACC_CLAUSE_VECTOR; + else if (!strcmp ("vector_length", p)) result = PRAGMA_OACC_CLAUSE_VECTOR_LENGTH; else if (flag_cilkplus && !strcmp ("vectorlength", p)) result = PRAGMA_CILK_CLAUSE_VECTORLENGTH; @@ -27691,11 +27807,13 @@ cp_parser_omp_clause_name (cp_parser *parser) case 'w': if (!strcmp ("wait", p)) result = PRAGMA_OACC_CLAUSE_WAIT; + else if (!strcmp ("worker", p)) + result = PRAGMA_OACC_CLAUSE_WORKER; break; } } - if (result != PRAGMA_OMP_CLAUSE_NONE) + if (consume_token && result != PRAGMA_OMP_CLAUSE_NONE) cp_lexer_consume_token (parser->lexer); return result; @@ -27893,6 +28011,8 @@ cp_parser_omp_var_list (cp_parser *parser, enum omp_clause_code kind, tree list) copyout ( variable-list ) create ( variable-list ) delete ( variable-list ) + device_resident ( variable-list ) + link ( variable-list ) present ( variable-list ) present_or_copy ( variable-list ) pcopy ( variable-list ) @@ -27928,10 +28048,15 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind, case PRAGMA_OACC_CLAUSE_DEVICE: kind = GOMP_MAP_FORCE_TO; break; + case PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT: + kind = GOMP_MAP_DEVICE_RESIDENT; + break; case PRAGMA_OACC_CLAUSE_HOST: - case PRAGMA_OACC_CLAUSE_SELF: kind = GOMP_MAP_FORCE_FROM; break; + case PRAGMA_OACC_CLAUSE_LINK: + kind = GOMP_MAP_LINK; + break; case PRAGMA_OACC_CLAUSE_PRESENT: kind = GOMP_MAP_FORCE_PRESENT; break; @@ -27999,43 +28124,360 @@ cp_parser_oacc_data_clause_deviceptr (cp_parser *parser, tree list) return list; } +/* Attempt to statically determine when the number T isn't positive. + Warn if we determined this and return positive one as the new + expression. */ +static tree +require_positive_expr (tree t, location_t loc, const char *str) +{ + tree c = fold_build2_loc (loc, LE_EXPR, boolean_type_node, t, + build_int_cst (TREE_TYPE (t), 0)); + if (c == boolean_true_node) + { + warning_at (loc, 0, + "%<%s%> value must be positive", str); + t = integer_one_node; + } + return t; +} + /* OpenACC: - vector_length ( expression ) */ + num_gangs ( expression ) + num_workers ( expression ) + vector_length ( expression ) + + OpenMP 2.5: + num_threads ( expression ) */ static tree -cp_parser_oacc_clause_vector_length (cp_parser *parser, tree list) +cp_parser_omp_positive_int_clause (cp_parser *parser, pragma_omp_clause c_kind, + const char *str, tree list) { - tree t, c; - location_t location = cp_lexer_peek_token (parser->lexer)->location; - bool error = false; + omp_clause_code kind; + switch (c_kind) + { + default: + gcc_unreachable (); + case PRAGMA_OACC_CLAUSE_NUM_GANGS: + kind = OMP_CLAUSE_NUM_GANGS; + break; + case PRAGMA_OMP_CLAUSE_NUM_THREADS: + kind = OMP_CLAUSE_NUM_THREADS; + break; + case PRAGMA_OACC_CLAUSE_NUM_WORKERS: + kind = OMP_CLAUSE_NUM_WORKERS; + break; + case PRAGMA_OACC_CLAUSE_VECTOR_LENGTH: + kind = OMP_CLAUSE_VECTOR_LENGTH; + break; + } + + location_t loc = cp_lexer_peek_token (parser->lexer)->location; if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN)) return list; - t = cp_parser_condition (parser); - if (t == error_mark_node || !INTEGRAL_TYPE_P (TREE_TYPE (t))) - { - error_at (location, "expected positive integer expression"); - error = true; - } + tree t = cp_parser_assignment_expression (parser, NULL, false, false); - if (error || !cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN)) - { - cp_parser_skip_to_closing_parenthesis (parser, /*recovering=*/true, + if (t == error_mark_node + || !cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN)) + cp_parser_skip_to_closing_parenthesis (parser, /*recovering=*/true, /*or_comma=*/false, /*consume_paren=*/true); + + check_no_duplicate_clause (list, kind, str, loc); + + tree c = build_omp_clause (loc, kind); + OMP_CLAUSE_OPERAND (c, 0) = t; + OMP_CLAUSE_CHAIN (c) = list; + return c; +} + +/* OpenACC: + gang [( gang_expr_list )] + worker [( expression )] + vector [( expression )] */ + +static tree +cp_parser_oacc_shape_clause (cp_parser *parser, pragma_omp_clause c_kind, + const char *str, tree list) +{ + omp_clause_code kind; + const char *id = "num"; + cp_lexer *lexer = parser->lexer; + + switch (c_kind) + { + default: + gcc_unreachable (); + case PRAGMA_OACC_CLAUSE_GANG: + kind = OMP_CLAUSE_GANG; + break; + case PRAGMA_OACC_CLAUSE_VECTOR: + kind = OMP_CLAUSE_VECTOR; + id = "length"; + break; + case PRAGMA_OACC_CLAUSE_WORKER: + kind = OMP_CLAUSE_WORKER; + break; + } + + tree op0 = NULL_TREE, op1 = NULL_TREE; + location_t loc = cp_lexer_peek_token (lexer)->location; + + if (cp_lexer_next_token_is (lexer, CPP_OPEN_PAREN)) + { + tree *op_to_parse = &op0; + cp_lexer_consume_token (lexer); + + do + { + if (cp_lexer_next_token_is (lexer, CPP_NAME) + || cp_lexer_next_token_is (lexer, CPP_KEYWORD)) + { + tree name_kind = cp_lexer_peek_token (lexer)->u.value; + const char *p = IDENTIFIER_POINTER (name_kind); + if (kind == OMP_CLAUSE_GANG && strcmp ("static", p) == 0) + { + cp_lexer_consume_token (lexer); + if (!cp_parser_require (parser, CPP_COLON, RT_COLON)) + { + cp_parser_skip_to_closing_parenthesis (parser, false, + false, true); + return list; + } + op_to_parse = &op1; + if (cp_lexer_next_token_is (lexer, CPP_MULT)) + { + if (*op_to_parse != NULL_TREE) + { + cp_parser_error (parser, + "duplicate % argument"); + cp_parser_skip_to_closing_parenthesis (parser, + false, false, + true); + return list; + } + cp_lexer_consume_token (lexer); + *op_to_parse = integer_minus_one_node; + if (cp_lexer_next_token_is (lexer, CPP_COMMA)) + cp_lexer_consume_token (lexer); + continue; + } + } + else if (strcmp (id, p) == 0) + { + op_to_parse = &op0; + cp_lexer_consume_token (lexer); + if (!cp_parser_require (parser, CPP_COLON, RT_COLON)) + { + cp_parser_skip_to_closing_parenthesis (parser, false, + false, true); + return list; + } + } + else + { + if (kind == OMP_CLAUSE_GANG) + cp_parser_error (parser, + "expected %<%num%> or %"); + else if (kind == OMP_CLAUSE_VECTOR) + cp_parser_error (parser, "expected %"); + else + cp_parser_error (parser, "expected %"); + cp_parser_skip_to_closing_parenthesis (parser, false, false, + true); + return list; + } + } + + if (*op_to_parse != NULL_TREE) + { + cp_parser_error (parser, "duplicate operand to clause"); + cp_parser_skip_to_closing_parenthesis (parser, false, false, + true); + return list; + } + + location_t expr_loc = cp_lexer_peek_token (lexer)->location; + tree expr = cp_parser_assignment_expression (parser, NULL, false, + false); + if (expr == error_mark_node) + { + cp_parser_skip_to_closing_parenthesis (parser, false, false, + true); + return list; + } + + mark_exp_read (expr); + require_positive_expr (expr, expr_loc, str); + *op_to_parse = expr; + + if (cp_lexer_next_token_is (lexer, CPP_COMMA)) + cp_lexer_consume_token (lexer); + } + while (!cp_lexer_next_token_is (lexer, CPP_CLOSE_PAREN)); + cp_lexer_consume_token (lexer); + } + + check_no_duplicate_clause (list, kind, str, loc); + + tree c = build_omp_clause (loc, kind); + if (op0) + OMP_CLAUSE_OPERAND (c, 0) = op0; + if (op1) + OMP_CLAUSE_OPERAND (c, 1) = op1; + OMP_CLAUSE_CHAIN (c) = list; + return c; +} + +/* OpenACC 2.0: + device_type ( size-expr-list ) clauses */ + +static tree +cp_parser_oacc_clause_device_type (cp_parser *parser, omp_clause_mask mask, + tree list, cp_token *pragma_tok) +{ + tree c, clauses; + location_t loc; + int dev_id = GOMP_DEVICE_NONE; + + loc = cp_lexer_peek_token (parser->lexer)->location; + if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN)) + return list; + + if (cp_lexer_next_token_is (parser->lexer, CPP_MULT)) + { + cp_lexer_consume_token (parser->lexer); + dev_id = GOMP_DEVICE_DEFAULT; + if (!cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN)) + return list; + } + else + { + do + { + tree keyword = error_mark_node; + int dev = 0; + + if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)) + { + keyword = cp_lexer_peek_token (parser->lexer)->u.value; + cp_lexer_consume_token (parser->lexer); + } + + if (keyword == error_mark_node) + { + error_at (loc, "expected keyword or %<)%>"); + cp_parser_skip_to_closing_parenthesis (parser, true, false, + true); + return list; + } + + dev = oacc_extract_device_id (IDENTIFIER_POINTER (keyword)); + if (dev) + dev_id |= 1 << dev; + + if (cp_lexer_next_token_is (parser->lexer, CPP_COMMA)) + cp_lexer_consume_token (parser->lexer); + } + while (cp_lexer_next_token_is_not (parser->lexer, CPP_CLOSE_PAREN)); + + /* Consume the trailing ')'. */ + cp_lexer_consume_token (parser->lexer); + } + + c = build_omp_clause (loc, OMP_CLAUSE_DEVICE_TYPE); + clauses = cp_parser_oacc_all_clauses (parser, mask, "device_type", + pragma_tok, 0, false, false); + OMP_CLAUSE_DEVICE_TYPE_CLAUSES (c) = clauses; + OMP_CLAUSE_DEVICE_TYPE_DEVICES (c) = build_int_cst (integer_type_node, + dev_id); + OMP_CLAUSE_CHAIN (c) = list; + return c; +} + +/* OpenACC 2.0: + tile ( size-expr-list ) */ + +static tree +cp_parser_oacc_clause_tile (cp_parser *parser, tree list, location_t here) +{ + tree c, num = error_mark_node; + HOST_WIDE_INT n; + location_t loc; + tree tile = NULL_TREE; + vec *tvec = make_tree_vector (); + + check_no_duplicate_clause (list, OMP_CLAUSE_TILE, "tile", here); + + loc = cp_lexer_peek_token (parser->lexer)->location; + if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN)) + { + release_tree_vector (tvec); return list; } - check_no_duplicate_clause (list, OMP_CLAUSE_VECTOR_LENGTH, "vector_length", - location); + do + { + if (cp_lexer_next_token_is (parser->lexer, CPP_MULT)) + { + cp_lexer_consume_token (parser->lexer); + num = integer_minus_one_node; + } + else + { + bool non_constant = false; + num = cp_parser_constant_expression (parser, true, &non_constant); - c = build_omp_clause (location, OMP_CLAUSE_VECTOR_LENGTH); - OMP_CLAUSE_VECTOR_LENGTH_EXPR (c) = t; + if (num == error_mark_node) + { + cp_parser_skip_to_closing_parenthesis (parser, true, false, + true); + release_tree_vector (tvec); + return list; + } + + num = fold_non_dependent_expr (num); + + if (non_constant + || !INTEGRAL_TYPE_P (TREE_TYPE (num)) + || !tree_fits_shwi_p (num) + || (n = tree_to_shwi (num)) <= 0 + || (int) n != n) + { + error_at (loc, + "tile argument needs positive constant integer " + "expression"); + release_tree_vector (tvec); + cp_parser_skip_to_closing_parenthesis (parser, true, false, + true); + return list; + } + } + + if (num == error_mark_node) + { + error_at (loc, "expected positive integer or %<)%>"); + release_tree_vector (tvec); + return list; + } + + vec_safe_push (tvec, num); + if (cp_lexer_next_token_is (parser->lexer, CPP_COMMA)) + cp_lexer_consume_token (parser->lexer); + } + while (cp_lexer_next_token_is_not (parser->lexer, CPP_CLOSE_PAREN)); + + /* Consume the trailing ')'. */ + cp_lexer_consume_token (parser->lexer); + + c = build_omp_clause (loc, OMP_CLAUSE_TILE); + tile = build_tree_list_vec (tvec); + OMP_CLAUSE_TILE_LIST (c) = tile; OMP_CLAUSE_CHAIN (c) = list; - list = c; - - return list; + release_tree_vector (tvec); + return c; } /* OpenACC 2.0 @@ -28054,7 +28496,8 @@ cp_parser_oacc_wait_list (cp_parser *parser, location_t clause_loc, tree list) if (args == NULL || args->length () == 0) { - cp_parser_error (parser, "expected integer expression before ')'"); + cp_parser_error (parser, + "expected integer expression list before %<)%>"); if (args != NULL) release_tree_vector (args); return list; @@ -28148,7 +28591,8 @@ cp_parser_omp_clause_collapse (cp_parser *parser, tree list, location_t location default ( shared | none ) */ static tree -cp_parser_omp_clause_default (cp_parser *parser, tree list, location_t location) +cp_parser_omp_clause_default (cp_parser *parser, tree list, + location_t location, bool is_omp) { enum omp_clause_default_kind kind = OMP_CLAUSE_DEFAULT_UNSPECIFIED; tree c; @@ -28169,7 +28613,7 @@ cp_parser_omp_clause_default (cp_parser *parser, tree list, location_t location) break; case 's': - if (strcmp ("shared", p) != 0) + if (strcmp ("shared", p) != 0 || !is_omp) goto invalid_kind; kind = OMP_CLAUSE_DEFAULT_SHARED; break; @@ -28183,7 +28627,10 @@ cp_parser_omp_clause_default (cp_parser *parser, tree list, location_t location) else { invalid_kind: - cp_parser_error (parser, "expected % or %"); + if (is_omp) + cp_parser_error (parser, "expected % or %"); + else + cp_parser_error (parser, "expected %"); } if (!cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN)) @@ -28291,109 +28738,6 @@ cp_parser_omp_clause_nowait (cp_parser * /*parser*/, return c; } -/* OpenACC: - num_gangs ( expression ) */ - -static tree -cp_parser_omp_clause_num_gangs (cp_parser *parser, tree list) -{ - tree t, c; - location_t location = cp_lexer_peek_token (parser->lexer)->location; - - if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN)) - return list; - - t = cp_parser_condition (parser); - - if (t == error_mark_node - || !cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN)) - cp_parser_skip_to_closing_parenthesis (parser, /*recovering=*/true, - /*or_comma=*/false, - /*consume_paren=*/true); - - if (!INTEGRAL_TYPE_P (TREE_TYPE (t))) - { - error_at (location, "expected positive integer expression"); - return list; - } - - check_no_duplicate_clause (list, OMP_CLAUSE_NUM_GANGS, "num_gangs", location); - - c = build_omp_clause (location, 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 ) */ - -static tree -cp_parser_omp_clause_num_threads (cp_parser *parser, tree list, - location_t location) -{ - tree t, c; - - if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN)) - return list; - - t = cp_parser_expression (parser); - - if (t == error_mark_node - || !cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN)) - cp_parser_skip_to_closing_parenthesis (parser, /*recovering=*/true, - /*or_comma=*/false, - /*consume_paren=*/true); - - check_no_duplicate_clause (list, OMP_CLAUSE_NUM_THREADS, - "num_threads", location); - - c = build_omp_clause (location, OMP_CLAUSE_NUM_THREADS); - OMP_CLAUSE_NUM_THREADS_EXPR (c) = t; - OMP_CLAUSE_CHAIN (c) = list; - - return c; -} - -/* OpenACC: - num_workers ( expression ) */ - -static tree -cp_parser_omp_clause_num_workers (cp_parser *parser, tree list) -{ - tree t, c; - location_t location = cp_lexer_peek_token (parser->lexer)->location; - - if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN)) - return list; - - t = cp_parser_condition (parser); - - if (t == error_mark_node - || !cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN)) - cp_parser_skip_to_closing_parenthesis (parser, /*recovering=*/true, - /*or_comma=*/false, - /*consume_paren=*/true); - - if (!INTEGRAL_TYPE_P (TREE_TYPE (t))) - { - error_at (location, "expected positive integer expression"); - return list; - } - - check_no_duplicate_clause (list, OMP_CLAUSE_NUM_WORKERS, "num_gangs", - location); - - c = build_omp_clause (location, OMP_CLAUSE_NUM_WORKERS); - OMP_CLAUSE_NUM_WORKERS_EXPR (c) = t; - OMP_CLAUSE_CHAIN (c) = list; - list = c; - - return list; -} - /* OpenMP 2.5: ordered */ @@ -28613,27 +28957,20 @@ cp_parser_omp_clause_schedule (cp_parser *parser, tree list, location_t location } /* OpenMP 3.0: - untied */ + untied -static tree -cp_parser_omp_clause_untied (cp_parser * /*parser*/, - tree list, location_t location) -{ - tree c; - - check_no_duplicate_clause (list, OMP_CLAUSE_UNTIED, "untied", location); - - c = build_omp_clause (location, OMP_CLAUSE_UNTIED); - OMP_CLAUSE_CHAIN (c) = list; - return c; -} - -/* OpenMP 4.0: + OpenMP 4.0: inbranch - notinbranch */ + notinbranch + + OpenACC 2.0: + auto + independent + nohost + seq */ static tree -cp_parser_omp_clause_branch (cp_parser * /*parser*/, enum omp_clause_code code, +cp_parser_omp_simple_clause (cp_parser * /*parser*/, enum omp_clause_code code, tree list, location_t location) { check_no_duplicate_clause (list, code, omp_clause_code_name[code], location); @@ -29121,16 +29458,66 @@ cp_parser_oacc_clause_async (cp_parser *parser, tree list) return list; } +/* OpenACC 2.0: + bind ( identifier ) + bind ( string-literal ) */ + +static tree +cp_parser_oacc_clause_bind (cp_parser *parser, tree list) +{ + location_t loc = cp_lexer_peek_token (parser->lexer)->location; + bool save_translate_strings_p = parser->translate_strings_p; + + parser->translate_strings_p = false; + if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN)) + { + parser->translate_strings_p = save_translate_strings_p; + return list; + } + if (cp_lexer_next_token_is (parser->lexer, CPP_NAME) + || cp_lexer_next_token_is (parser->lexer, CPP_STRING)) + { + tree t; + + if (cp_lexer_peek_token (parser->lexer)->type == CPP_STRING) + { + t = cp_lexer_peek_token (parser->lexer)->u.value; + cp_lexer_consume_token (parser->lexer); + } + else + t = cp_parser_id_expression (parser, /*template_p=*/false, + /*check_dependency_p=*/true, + /*template_p=*/NULL, + /*declarator_p=*/false, + /*optional_p=*/false); + if (t == error_mark_node) + return t; + + tree c = build_omp_clause (loc, OMP_CLAUSE_BIND); + OMP_CLAUSE_BIND_NAME (c) = t; + OMP_CLAUSE_CHAIN (c) = list; + list = c; + } + else + cp_parser_error (parser, "expected identifier or character string literal"); + parser->translate_strings_p = save_translate_strings_p; + cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN); + return list; +} + /* Parse all OpenACC clauses. The set clauses allowed by the directive - is a bitmask in MASK. Return the list of clauses found. */ + is a bitmask in MASK. DTYPE_MASK denotes clauses which may follow a + device_type mask. Return the list of clauses found. */ -static tree +tree cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, - const char *where, cp_token *pragma_tok, - bool finish_p = true) + const char *where, cp_token *pragma_tok, + omp_clause_mask dtype_mask = 0, + bool finish_p = true, bool scan_dtype = true) { tree clauses = NULL; bool first = true; + bool seen_dtype = false; while (cp_lexer_next_token_is_not (parser->lexer, CPP_PRAGMA_EOL)) { @@ -29142,15 +29529,35 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, if (!first && cp_lexer_next_token_is (parser->lexer, CPP_COMMA)) cp_lexer_consume_token (parser->lexer); + if (!scan_dtype && cp_parser_omp_clause_name (parser, false) + == PRAGMA_OACC_CLAUSE_DEVICE_TYPE) + return clauses; + here = cp_lexer_peek_token (parser->lexer)->location; c_kind = cp_parser_omp_clause_name (parser); + if (seen_dtype && c_kind != PRAGMA_OMP_CLAUSE_NONE + && c_kind != PRAGMA_OACC_CLAUSE_DEVICE_TYPE) + { + error_at (here, "invalid clauses following device_type"); + goto saw_error; + } + switch (c_kind) { case PRAGMA_OACC_CLAUSE_ASYNC: clauses = cp_parser_oacc_clause_async (parser, clauses); c_name = "async"; break; + case PRAGMA_OACC_CLAUSE_AUTO: + clauses = cp_parser_omp_simple_clause (parser, OMP_CLAUSE_AUTO, + clauses, here); + c_name = "auto"; + break; + case PRAGMA_OACC_CLAUSE_BIND: + clauses = cp_parser_oacc_clause_bind (parser, clauses); + c_name = "bind"; + break; case PRAGMA_OACC_CLAUSE_COLLAPSE: clauses = cp_parser_omp_clause_collapse (parser, clauses, here); c_name = "collapse"; @@ -29175,29 +29582,66 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); c_name = "delete"; break; + case PRAGMA_OMP_CLAUSE_DEFAULT: + clauses = cp_parser_omp_clause_default (parser, clauses, here, + false); + c_name = "default"; + break; case PRAGMA_OACC_CLAUSE_DEVICE: clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); c_name = "device"; break; + case PRAGMA_OACC_CLAUSE_DEVICE_TYPE: + clauses = cp_parser_oacc_clause_device_type (parser, dtype_mask, + clauses, pragma_tok); + c_name = "device_type"; + seen_dtype = true; + break; + case PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT: + clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "device_resident"; + break; case PRAGMA_OACC_CLAUSE_DEVICEPTR: clauses = cp_parser_oacc_data_clause_deviceptr (parser, clauses); c_name = "deviceptr"; break; - case PRAGMA_OACC_CLAUSE_HOST: - clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); - c_name = "host"; - break; case PRAGMA_OACC_CLAUSE_IF: clauses = cp_parser_omp_clause_if (parser, clauses, here); c_name = "if"; break; + case PRAGMA_OACC_CLAUSE_INDEPENDENT: + clauses = cp_parser_omp_simple_clause (parser, + OMP_CLAUSE_INDEPENDENT, + clauses, here); + c_name = "independent"; + break; + case PRAGMA_OACC_CLAUSE_GANG: + c_name = "gang"; + clauses = cp_parser_oacc_shape_clause (parser, c_kind, c_name, + clauses); + break; + case PRAGMA_OACC_CLAUSE_HOST: + clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "host"; + break; + case PRAGMA_OACC_CLAUSE_LINK: + clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "link"; + break; + case PRAGMA_OACC_CLAUSE_NOHOST: + clauses = cp_parser_omp_simple_clause (parser, OMP_CLAUSE_NOHOST, + clauses, here); + c_name = "nohost"; + break; case PRAGMA_OACC_CLAUSE_NUM_GANGS: - clauses = cp_parser_omp_clause_num_gangs (parser, clauses); c_name = "num_gangs"; + clauses = cp_parser_omp_positive_int_clause (parser, c_kind, c_name, + clauses); break; case PRAGMA_OACC_CLAUSE_NUM_WORKERS: - clauses = cp_parser_omp_clause_num_workers (parser, clauses); c_name = "num_workers"; + clauses = cp_parser_omp_positive_int_clause (parser, c_kind, c_name, + clauses); break; case PRAGMA_OACC_CLAUSE_PRESENT: clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); @@ -29219,22 +29663,48 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); c_name = "present_or_create"; break; + case PRAGMA_OACC_CLAUSE_PRIVATE: + clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_PRIVATE, + clauses); + c_name = "private"; + break; case PRAGMA_OACC_CLAUSE_REDUCTION: clauses = cp_parser_omp_clause_reduction (parser, clauses); c_name = "reduction"; break; - case PRAGMA_OACC_CLAUSE_SELF: - clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); - c_name = "self"; + case PRAGMA_OACC_CLAUSE_SEQ: + clauses = cp_parser_omp_simple_clause (parser, OMP_CLAUSE_SEQ, + clauses, here); + c_name = "seq"; + break; + case PRAGMA_OACC_CLAUSE_TILE: + clauses = cp_parser_oacc_clause_tile (parser, clauses, here); + c_name = "tile"; + break; + case PRAGMA_OACC_CLAUSE_USE_DEVICE: + clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_USE_DEVICE, + clauses); + c_name = "use_device"; + break; + case PRAGMA_OACC_CLAUSE_VECTOR: + c_name = "vector"; + clauses = cp_parser_oacc_shape_clause (parser, c_kind, c_name, + clauses); break; case PRAGMA_OACC_CLAUSE_VECTOR_LENGTH: - clauses = cp_parser_oacc_clause_vector_length (parser, clauses); c_name = "vector_length"; + clauses = cp_parser_omp_positive_int_clause (parser, c_kind, c_name, + clauses); break; case PRAGMA_OACC_CLAUSE_WAIT: clauses = cp_parser_oacc_clause_wait (parser, clauses); c_name = "wait"; break; + case PRAGMA_OACC_CLAUSE_WORKER: + c_name = "worker"; + clauses = cp_parser_oacc_shape_clause (parser, c_kind, c_name, + clauses); + break; default: cp_parser_error (parser, "expected %<#pragma acc%> clause"); goto saw_error; @@ -29251,11 +29721,17 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, } } + if (!scan_dtype) + return clauses; + saw_error: cp_parser_skip_to_pragma_eol (parser, pragma_tok); if (finish_p) - return finish_omp_clauses (clauses); + { + clauses = oacc_filter_device_types (clauses); + return finish_omp_clauses (clauses, true); + } return clauses; } @@ -29304,7 +29780,7 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask, break; case PRAGMA_OMP_CLAUSE_DEFAULT: clauses = cp_parser_omp_clause_default (parser, clauses, - token->location); + token->location, true); c_name = "default"; break; case PRAGMA_OMP_CLAUSE_FINAL: @@ -29335,9 +29811,9 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask, c_name = "nowait"; break; case PRAGMA_OMP_CLAUSE_NUM_THREADS: - clauses = cp_parser_omp_clause_num_threads (parser, clauses, - token->location); c_name = "num_threads"; + clauses = cp_parser_omp_positive_int_clause (parser, c_kind, c_name, + clauses); break; case PRAGMA_OMP_CLAUSE_ORDERED: clauses = cp_parser_omp_clause_ordered (parser, clauses, @@ -29364,19 +29840,19 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask, c_name = "shared"; break; case PRAGMA_OMP_CLAUSE_UNTIED: - clauses = cp_parser_omp_clause_untied (parser, clauses, - token->location); + clauses = cp_parser_omp_simple_clause (parser, OMP_CLAUSE_UNTIED, + clauses, token->location); c_name = "untied"; break; case PRAGMA_OMP_CLAUSE_INBRANCH: case PRAGMA_CILK_CLAUSE_MASK: - clauses = cp_parser_omp_clause_branch (parser, OMP_CLAUSE_INBRANCH, + clauses = cp_parser_omp_simple_clause (parser, OMP_CLAUSE_INBRANCH, clauses, token->location); c_name = "inbranch"; break; case PRAGMA_OMP_CLAUSE_NOTINBRANCH: case PRAGMA_CILK_CLAUSE_NOMASK: - clauses = cp_parser_omp_clause_branch (parser, + clauses = cp_parser_omp_simple_clause (parser, OMP_CLAUSE_NOTINBRANCH, clauses, token->location); c_name = "notinbranch"; @@ -29507,7 +29983,7 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask, if (!(flag_cilkplus && pragma_tok == NULL)) cp_parser_skip_to_pragma_eol (parser, pragma_tok); if (finish_p) - return finish_omp_clauses (clauses); + return finish_omp_clauses (clauses, false); return clauses; } @@ -30501,7 +30977,7 @@ cp_parser_omp_for_loop (cp_parser *parser, enum tree_code code, tree clauses, { c = build_omp_clause (loc, OMP_CLAUSE_PRIVATE); OMP_CLAUSE_DECL (c) = decl; - c = finish_omp_clauses (c); + c = finish_omp_clauses (c, false); if (c) { OMP_CLAUSE_CHAIN (c) = clauses; @@ -30640,7 +31116,7 @@ cp_omp_split_clauses (location_t loc, enum tree_code code, c_omp_split_clauses (loc, code, mask, clauses, cclauses); for (i = 0; i < C_OMP_CLAUSE_SPLIT_COUNT; i++) if (cclauses[i]) - cclauses[i] = finish_omp_clauses (cclauses[i]); + cclauses[i] = finish_omp_clauses (cclauses[i], false); } /* OpenMP 4.0: @@ -31490,7 +31966,7 @@ cp_parser_oacc_cache (cp_parser *parser, cp_token *pragma_tok) tree stmt, clauses; clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE__CACHE_, NULL_TREE); - clauses = finish_omp_clauses (clauses); + clauses = finish_omp_clauses (clauses, true); cp_parser_require_pragma_eol (parser, cp_lexer_peek_token (parser->lexer)); @@ -31537,14 +32013,36 @@ cp_parser_oacc_data (cp_parser *parser, cp_token *pragma_tok) return stmt; } +#define OACC_HOST_DATA_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_USE_DEVICE) ) + +/* OpenACC 2.0: + # pragma acc host_data new-line + structured-block */ + +static tree +cp_parser_oacc_host_data (cp_parser *parser, cp_token *pragma_tok) +{ + tree stmt, clauses, block; + unsigned int save; + + clauses = cp_parser_oacc_all_clauses (parser, OACC_HOST_DATA_CLAUSE_MASK, + "#pragma acc host_data", pragma_tok); + + block = begin_omp_parallel (); + save = cp_parser_begin_omp_structured_block (parser); + cp_parser_statement (parser, NULL_TREE, false, NULL); + cp_parser_end_omp_structured_block (parser, save); + stmt = finish_oacc_host_data (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 \ @@ -31567,23 +32065,18 @@ static tree cp_parser_oacc_enter_exit_data (cp_parser *parser, cp_token *pragma_tok, bool enter) { + location_t loc = pragma_tok->location; tree stmt, clauses; + const char *p = ""; - if (cp_lexer_next_token_is (parser->lexer, CPP_PRAGMA_EOL) - || cp_lexer_next_token_is_not (parser->lexer, CPP_NAME)) - { - cp_parser_error (parser, enter - ? "expected % in %<#pragma acc enter data%>" - : "expected % in %<#pragma acc exit data%>"); - cp_parser_skip_to_pragma_eol (parser, pragma_tok); - return NULL_TREE; - } + if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)) + p = IDENTIFIER_POINTER (cp_lexer_peek_token (parser->lexer)->u.value); - const char *p = - IDENTIFIER_POINTER (cp_lexer_peek_token (parser->lexer)->u.value); if (strcmp (p, "data") != 0) { - cp_parser_error (parser, "invalid pragma"); + error_at (loc, enter + ? "expected % after %<#pragma acc enter%>" + : "expected % after %<#pragma acc exit%>"); cp_parser_skip_to_pragma_eol (parser, pragma_tok); return NULL_TREE; } @@ -31599,53 +32092,68 @@ cp_parser_oacc_enter_exit_data (cp_parser *parser, cp_token *pragma_tok, if (find_omp_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE) { - error_at (pragma_tok->location, - "%<#pragma acc enter data%> has no data movement clause"); + error_at (loc, "%<#pragma acc %s data%> has no data movement clause", + enter ? "enter" : "exit"); return NULL_TREE; } stmt = enter ? make_node (OACC_ENTER_DATA) : make_node (OACC_EXIT_DATA); TREE_TYPE (stmt) = void_type_node; OMP_STANDALONE_CLAUSES (stmt) = clauses; - SET_EXPR_LOCATION (stmt, pragma_tok->location); + SET_EXPR_LOCATION (stmt, loc); add_stmt (stmt); 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_OACC_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_OACC_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)) +/* Split the 'clauses' into a set of 'loop' clauses and a set of + 'not-loop' clauses. */ static tree -cp_parser_oacc_kernels (cp_parser *parser, cp_token *pragma_tok) +oacc_split_loop_clauses (tree clauses, tree *not_loop_clauses) { - tree stmt, clauses, block; - unsigned int save; - - clauses = cp_parser_oacc_all_clauses (parser, OACC_KERNELS_CLAUSE_MASK, - "#pragma acc kernels", pragma_tok); - - block = begin_omp_parallel (); - save = cp_parser_begin_omp_structured_block (parser); - cp_parser_statement (parser, NULL_TREE, false, NULL); - cp_parser_end_omp_structured_block (parser, save); - stmt = finish_oacc_kernels (clauses, block); - return stmt; + tree loop_clauses, next, c; + + loop_clauses = *not_loop_clauses = NULL_TREE; + + for (; clauses ; clauses = next) + { + next = OMP_CLAUSE_CHAIN (clauses); + + switch (OMP_CLAUSE_CODE (clauses)) + { + case OMP_CLAUSE_COLLAPSE: + case OMP_CLAUSE_REDUCTION: + case OMP_CLAUSE_GANG: + case OMP_CLAUSE_VECTOR: + case OMP_CLAUSE_WORKER: + case OMP_CLAUSE_AUTO: + case OMP_CLAUSE_SEQ: + OMP_CLAUSE_CHAIN (clauses) = loop_clauses; + loop_clauses = clauses; + break; + + case OMP_CLAUSE_PRIVATE: + c = build_omp_clause (OMP_CLAUSE_LOCATION (clauses), + OMP_CLAUSE_CODE (clauses)); + OMP_CLAUSE_DECL (c) = OMP_CLAUSE_DECL (clauses); + OMP_CLAUSE_CHAIN (c) = loop_clauses; + loop_clauses = c; + /* FALL THROUGH */ + + default: + OMP_CLAUSE_CHAIN (clauses) = *not_loop_clauses; + *not_loop_clauses = clauses; + break; + } + } + + if (*not_loop_clauses) + finish_omp_clauses (*not_loop_clauses, true); + + if (loop_clauses) + finish_omp_clauses (loop_clauses, true); + + return loop_clauses; } /* OpenACC 2.0: @@ -31654,16 +32162,43 @@ cp_parser_oacc_kernels (cp_parser *parser, cp_token *pragma_tok) #define OACC_LOOP_CLAUSE_MASK \ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COLLAPSE) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION)) + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_TYPE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_GANG) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WORKER) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_AUTO) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_INDEPENDENT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_TILE)) + +#define OACC_LOOP_CLAUSE_DEVICE_TYPE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COLLAPSE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_GANG) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WORKER) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_AUTO) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_INDEPENDENT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_TILE) ) static tree -cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok) +cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name, + omp_clause_mask mask, tree *cclauses) { tree stmt, clauses, block; int save; - clauses = cp_parser_oacc_all_clauses (parser, OACC_LOOP_CLAUSE_MASK, - "#pragma acc loop", pragma_tok); + strcat (p_name, " loop"); + mask |= OACC_LOOP_CLAUSE_MASK; + + clauses = cp_parser_oacc_all_clauses (parser, mask, p_name, pragma_tok, + OACC_LOOP_CLAUSE_DEVICE_TYPE_MASK, + cclauses == NULL); + + if (cclauses) + clauses = oacc_split_loop_clauses (clauses, cclauses); block = begin_omp_structured_block (); save = cp_parser_begin_omp_structured_block (parser); @@ -31674,6 +32209,31 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok) } /* 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_OACC_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_DEFAULT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_TYPE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_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)) + +#define OACC_KERNELS_CLAUSE_DEVICE_TYPE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) + +/* OpenACC 2.0: # pragma acc parallel oacc-parallel-clause[optseq] new-line structured-block */ @@ -31683,7 +32243,10 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok) | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_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_DEFAULT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_TYPE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_GANG) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS) \ @@ -31692,24 +32255,68 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok) | (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_PRIVATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT)) +#define OACC_PARALLEL_CLAUSE_DEVICE_TYPE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \ + | (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_VECTOR_LENGTH) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) + static tree -cp_parser_oacc_parallel (cp_parser *parser, cp_token *pragma_tok) +cp_parser_oacc_parallel_kernels (cp_parser *parser, cp_token *pragma_tok, + char *p_name, bool is_parallel) { tree stmt, clauses, block; unsigned int save; + cp_lexer *lexer = parser->lexer; + omp_clause_mask mask, dtype_mask; - clauses = cp_parser_oacc_all_clauses (parser, OACC_PARALLEL_CLAUSE_MASK, - "#pragma acc parallel", pragma_tok); + if (is_parallel) + { + mask = OACC_PARALLEL_CLAUSE_MASK; + strcat (p_name, " parallel"); + } + else + { + mask = OACC_KERNELS_CLAUSE_MASK; + strcat (p_name, " kernels"); + } + + if (cp_lexer_next_token_is (lexer, CPP_NAME)) + { + stmt = cp_lexer_peek_token (lexer)->u.value; + if (!strcmp ("loop", IDENTIFIER_POINTER (stmt))) + { + tree combined_clauses = NULL_TREE; + + cp_lexer_consume_token (lexer); + mask |= OACC_LOOP_CLAUSE_MASK; + block = begin_omp_parallel (); + cp_parser_oacc_loop (parser, pragma_tok, p_name, mask, + &combined_clauses); + stmt = is_parallel ? finish_oacc_parallel (combined_clauses, block) + : finish_oacc_kernels (combined_clauses, block); + return stmt; + } + } + + dtype_mask = is_parallel ? OACC_PARALLEL_CLAUSE_DEVICE_TYPE_MASK + : OACC_KERNELS_CLAUSE_DEVICE_TYPE_MASK; + + clauses = cp_parser_oacc_all_clauses (parser, mask, p_name, pragma_tok, + dtype_mask); block = begin_omp_parallel (); save = cp_parser_begin_omp_structured_block (parser); cp_parser_statement (parser, NULL_TREE, false, NULL); cp_parser_end_omp_structured_block (parser, save); - stmt = finish_oacc_parallel (clauses, block); + stmt = is_parallel ? finish_oacc_parallel (clauses, block) + : finish_oacc_kernels (clauses, block); return stmt; } @@ -31720,18 +32327,23 @@ cp_parser_oacc_parallel (cp_parser *parser, cp_token *pragma_tok) #define OACC_UPDATE_CLAUSE_MASK \ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_TYPE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_HOST) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SELF) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT)) +#define OACC_UPDATE_CLAUSE_DEVICE_TYPE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) + static tree cp_parser_oacc_update (cp_parser *parser, cp_token *pragma_tok) { tree stmt, clauses; clauses = cp_parser_oacc_all_clauses (parser, OACC_UPDATE_CLAUSE_MASK, - "#pragma acc update", pragma_tok); + "#pragma acc update", pragma_tok, + OACC_UPDATE_CLAUSE_DEVICE_TYPE_MASK); if (find_omp_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE) { @@ -31751,8 +32363,6 @@ cp_parser_oacc_update (cp_parser *parser, cp_token *pragma_tok) /* 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 \ @@ -31776,7 +32386,14 @@ cp_parser_oacc_wait (cp_parser *parser, cp_token *pragma_tok) } /* OpenMP 4.0: - # pragma omp declare simd declare-simd-clauses[optseq] new-line */ + # pragma omp declare simd declare-simd-clauses[optseq] new-line + + OpenACC 2.0a: + # pragma acc routine oacc-routine-clause[optseq] new-line + function-definition + + # pragma acc routine ( name ) oacc-routine-clause[optseq] new-line +*/ #define OMP_DECLARE_SIMD_CLAUSE_MASK \ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_SIMDLEN) \ @@ -31788,26 +32405,42 @@ cp_parser_oacc_wait (cp_parser *parser, cp_token *pragma_tok) static void cp_parser_omp_declare_simd (cp_parser *parser, cp_token *pragma_tok, - enum pragma_context context) + enum pragma_context context, bool is_omp) { - bool first_p = parser->omp_declare_simd == NULL; + bool first_p = is_omp ? parser->omp_declare_simd == NULL + : parser->oacc_routine == NULL; cp_omp_declare_simd_data data; if (first_p) { data.error_seen = false; data.fndecl_seen = false; data.tokens = vNULL; - parser->omp_declare_simd = &data; + if (is_omp) + parser->omp_declare_simd = &data; + else + parser->oacc_routine = &data; } while (cp_lexer_next_token_is_not (parser->lexer, CPP_PRAGMA_EOL) && cp_lexer_next_token_is_not (parser->lexer, CPP_EOF)) cp_lexer_consume_token (parser->lexer); + if (cp_lexer_next_token_is_not (parser->lexer, CPP_PRAGMA_EOL)) - parser->omp_declare_simd->error_seen = true; + { + if (is_omp) + parser->omp_declare_simd->error_seen = true; + else + parser->oacc_routine->error_seen = true; + } + cp_parser_require_pragma_eol (parser, pragma_tok); struct cp_token_cache *cp = cp_token_cache_new (pragma_tok, cp_lexer_peek_token (parser->lexer)); - parser->omp_declare_simd->tokens.safe_push (cp); + + if (is_omp) + parser->omp_declare_simd->tokens.safe_push (cp); + else + parser->oacc_routine->tokens.safe_push (cp); + if (first_p) { while (cp_lexer_next_token_is (parser->lexer, CPP_PRAGMA)) @@ -31827,14 +32460,23 @@ cp_parser_omp_declare_simd (cp_parser *parser, cp_token *pragma_tok, cp_parser_declaration_statement (parser); break; } - if (parser->omp_declare_simd + if (is_omp && parser->omp_declare_simd && !parser->omp_declare_simd->error_seen && !parser->omp_declare_simd->fndecl_seen) error_at (pragma_tok->location, "%<#pragma omp declare simd%> not immediately followed by " "function declaration or definition"); + else if (!is_omp && parser->oacc_routine + && !parser->oacc_routine->error_seen + && !parser->oacc_routine->fndecl_seen) + error_at (pragma_tok->location, + "%<#pragma acc routine%> not immediately followed by " + "function declaration or definition"); data.tokens.release (); - parser->omp_declare_simd = NULL; + if (is_omp) + parser->omp_declare_simd = NULL; + else + parser->oacc_routine = NULL; } } @@ -32410,7 +33052,7 @@ cp_parser_omp_declare (cp_parser *parser, cp_token *pragma_tok, { cp_lexer_consume_token (parser->lexer); cp_parser_omp_declare_simd (parser, pragma_tok, - context); + context, true); return; } cp_ensure_no_omp_declare_simd (parser); @@ -32438,6 +33080,167 @@ cp_parser_omp_declare (cp_parser *parser, cp_token *pragma_tok, cp_parser_require_pragma_eol (parser, pragma_tok); } +static void +cp_parser_oacc_routine_check_parallelism (tree clauses, location_t loc) +{ + /* Check of the presence if gang, worker, vector and seq clauses, and + throw an error if more than one of those clauses is specified. */ + int parallelism = 0; + tree c; + + for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) + switch (OMP_CLAUSE_CODE (c)) + { + case OMP_CLAUSE_GANG: + case OMP_CLAUSE_WORKER: + case OMP_CLAUSE_VECTOR: + case OMP_CLAUSE_SEQ: + ++parallelism; + break; + default: + break; + } + + if (parallelism > 1) + { + error_at (loc, "invalid combination of gang, worker, vector or seq for" + "%<#pragma acc routine%>"); + } +} + +/* OpenACC 2.0: + # pragma acc routine oacc-routine-clause[optseq] new-line + function-definition + + # pragma acc routine ( name ) oacc-routine-clause[optseq] new-line +*/ + +#define OACC_ROUTINE_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_TYPE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_GANG) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WORKER) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_BIND) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NOHOST)) + +#define OACC_ROUTINE_CLAUSE_DEVICE_TYPE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_GANG) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WORKER) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_BIND)) + +static void +cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok, + enum pragma_context context) +{ + tree name = NULL_TREE; + location_t here = cp_lexer_peek_token (parser->lexer)->location; + + //cp_lexer_consume_token (parser->lexer); + + /* Scan for optional '( name )'. */ + if (cp_lexer_next_token_is (parser->lexer, CPP_OPEN_PAREN)) + { + cp_lexer_consume_token (parser->lexer); + name = cp_parser_id_expression (parser, /*template_p=*/false, + /*check_dependency_p=*/true, + /*template_p=*/NULL, + /*declarator_p=*/false, + /*optional_p=*/false); + if (name == error_mark_node) + return; + + if (cp_lexer_next_token_is_not (parser->lexer, CPP_CLOSE_PAREN)) + { + error_at (cp_lexer_peek_token (parser->lexer)->location, + "expected %<)%>"); + return; + } + cp_lexer_consume_token (parser->lexer); + } + + /* If this routine construct doesn't explicitly have an optional 'name', + then handle it the same way as an omp declare simd. */ + if (!name) + { + cp_parser_omp_declare_simd (parser, pragma_tok, context, false); + cp_ensure_no_omp_declare_simd (parser); + return; + } + + /* Build a chain of clauses. */ + parser->lexer->in_pragma = true; + tree clauses = NULL_TREE; + clauses = cp_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK, + "#pragma acc routine", + cp_lexer_peek_token (parser->lexer), + OACC_ROUTINE_CLAUSE_DEVICE_TYPE_MASK); + + cp_parser_oacc_routine_check_parallelism (clauses, here); + + TREE_CHAIN (name) = clauses; + vec_safe_push (parser->named_oacc_routines, name); +} + +/* Finalize #pragma acc routine clauses after direct declarator has + been parsed, and put that into "omp declare target" attribute. */ + +static tree +cp_parser_late_parsing_oacc_routine (cp_parser *parser, tree attrs) +{ + struct cp_token_cache *ce; + cp_omp_declare_simd_data *data = parser->oacc_routine; + int i; + location_t here = UNKNOWN_LOCATION; + + if (!data->error_seen && data->fndecl_seen) + { + error ("%<#pragma acc routine%> not immediately followed by " + "a single function declaration or definition"); + data->error_seen = true; + return attrs; + } + if (data->error_seen) + return attrs; + + tree c, cl = NULL_TREE; + + FOR_EACH_VEC_ELT (data->tokens, i, ce) + { + cp_parser_push_lexer_for_tokens (parser, ce); + parser->lexer->in_pragma = true; + here = cp_lexer_peek_token (parser->lexer)->location; + gcc_assert (cp_lexer_peek_token (parser->lexer)->type == CPP_PRAGMA); + cp_token *pragma_tok = cp_lexer_consume_token (parser->lexer); + c = cp_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK, + "#pragma acc routine", pragma_tok); + cp_parser_pop_lexer (parser); + + if (cl == NULL_TREE) + cl = c; + else if (c != NULL_TREE) + { + OMP_CLAUSE_CHAIN (c) = cl; + cl = c; + TREE_CHAIN (c) = attrs; + if (processing_template_decl) + ATTR_IS_DEPENDENT (c) = 1; + attrs = c; + } + } + + cp_parser_oacc_routine_check_parallelism (cl, here); + + if (cl != NULL_TREE) + cl = tree_cons (NULL_TREE, cl, NULL_TREE); + + attrs = build_tree_list (get_identifier ("omp declare target"), cl); + data->fndecl_seen = true; + return attrs; +} + /* Main entry point to OpenMP statement pragmas. */ static void @@ -32449,6 +33252,9 @@ cp_parser_omp_construct (cp_parser *parser, cp_token *pragma_tok) switch (pragma_tok->pragma_kind) { + case PRAGMA_OACC_ATOMIC: + cp_parser_omp_atomic (parser, pragma_tok); + return; case PRAGMA_OACC_CACHE: stmt = cp_parser_oacc_cache (parser, pragma_tok); break; @@ -32461,14 +33267,22 @@ cp_parser_omp_construct (cp_parser *parser, cp_token *pragma_tok) case PRAGMA_OACC_EXIT_DATA: stmt = cp_parser_oacc_enter_exit_data (parser, pragma_tok, false); break; + case PRAGMA_OACC_HOST_DATA: + stmt = cp_parser_oacc_host_data (parser, pragma_tok); + break; case PRAGMA_OACC_KERNELS: - stmt = cp_parser_oacc_kernels (parser, pragma_tok); + strcpy (p_name, "#pragma acc"); + stmt = cp_parser_oacc_parallel_kernels (parser, pragma_tok, p_name, + false); break; case PRAGMA_OACC_LOOP: - stmt = cp_parser_oacc_loop (parser, pragma_tok); + strcpy (p_name, "#pragma acc"); + stmt = cp_parser_oacc_loop (parser, pragma_tok, p_name, mask, NULL); break; case PRAGMA_OACC_PARALLEL: - stmt = cp_parser_oacc_parallel (parser, pragma_tok); + strcpy (p_name, "#pragma acc"); + stmt = cp_parser_oacc_parallel_kernels (parser, pragma_tok, p_name, + true); break; case PRAGMA_OACC_UPDATE: stmt = cp_parser_oacc_update (parser, pragma_tok); @@ -32907,7 +33721,7 @@ cp_parser_pragma (cp_parser *parser, enum pragma_context context) parser->lexer->in_pragma = true; id = pragma_tok->pragma_kind; - if (id != PRAGMA_OMP_DECLARE_REDUCTION) + if (id != PRAGMA_OMP_DECLARE_REDUCTION && id != PRAGMA_OACC_ROUTINE) cp_ensure_no_omp_declare_simd (parser); switch (id) { @@ -33018,15 +33832,65 @@ cp_parser_pragma (cp_parser *parser, enum pragma_context context) cp_parser_omp_declare (parser, pragma_tok, context); return false; - case PRAGMA_OACC_CACHE: - case PRAGMA_OACC_DATA: case PRAGMA_OACC_ENTER_DATA: + if (context == pragma_stmt) + { + cp_parser_error (parser, "%<#pragma acc enter data%> may only be " + "used in compound statements"); + break; + } + else if (context != pragma_compound) + goto bad_stmt; + cp_parser_omp_construct (parser, pragma_tok); + return true; + case PRAGMA_OACC_EXIT_DATA: + if (context == pragma_stmt) + { + cp_parser_error (parser, "%<#pragma acc exit data%> may only be " + "used in compound statements"); + break; + } + else if (context != pragma_compound) + goto bad_stmt; + cp_parser_omp_construct (parser, pragma_tok); + return true; + + case PRAGMA_OACC_ROUTINE: + cp_parser_oacc_routine (parser, pragma_tok, context); + return false; + + case PRAGMA_OACC_UPDATE: + if (context == pragma_stmt) + { + cp_parser_error (parser, "%<#pragma acc update%> may only be " + "used in compound statements"); + break; + } + else if (context != pragma_compound) + goto bad_stmt; + cp_parser_omp_construct (parser, pragma_tok); + return true; + + case PRAGMA_OACC_WAIT: + if (context == pragma_stmt) + { + cp_parser_error (parser, "%<#pragma acc wait%> may only be " + "used in compound statements"); + break; + } + else if (context != pragma_compound) + goto bad_stmt; + cp_parser_omp_construct (parser, pragma_tok); + return true; + + case PRAGMA_OACC_ATOMIC: + case PRAGMA_OACC_CACHE: + case PRAGMA_OACC_DATA: + case PRAGMA_OACC_HOST_DATA: case PRAGMA_OACC_KERNELS: case PRAGMA_OACC_PARALLEL: case PRAGMA_OACC_LOOP: - case PRAGMA_OACC_UPDATE: - case PRAGMA_OACC_WAIT: case PRAGMA_OMP_ATOMIC: case PRAGMA_OMP_CRITICAL: case PRAGMA_OMP_DISTRIBUTE: @@ -33468,7 +34332,7 @@ cp_parser_cilk_for (cp_parser *parser, tree grain) tree clauses = build_omp_clause (EXPR_LOCATION (grain), OMP_CLAUSE_SCHEDULE); OMP_CLAUSE_SCHEDULE_KIND (clauses) = OMP_CLAUSE_SCHEDULE_CILKFOR; OMP_CLAUSE_SCHEDULE_CHUNK_EXPR (clauses) = grain; - clauses = finish_omp_clauses (clauses); + clauses = finish_omp_clauses (clauses, false); tree ret = cp_parser_omp_for_loop (parser, CILK_FOR, clauses, NULL); if (ret) diff --git gcc/cp/parser.h gcc/cp/parser.h index 76e5367..8eb5484 100644 --- gcc/cp/parser.h +++ gcc/cp/parser.h @@ -373,6 +373,10 @@ typedef struct GTY(()) cp_parser { necessary. */ cp_omp_declare_simd_data * GTY((skip)) cilk_simd_fn_info; + /* OpenACC specific parser information. */ + cp_omp_declare_simd_data * GTY((skip)) oacc_routine; + vec *named_oacc_routines; + /* Nonzero if parsing a parameter list where 'auto' should trigger an implicit template parameter. */ bool auto_is_implicit_function_template_parm_p; diff --git gcc/cp/pt.c gcc/cp/pt.c index 129517a..bbd54fe 100644 --- gcc/cp/pt.c +++ gcc/cp/pt.c @@ -8990,7 +8990,7 @@ apply_late_template_attributes (tree *decl_p, tree attributes, int attr_flags, clauses = tsubst_omp_clauses (clauses, true, args, complain, in_decl); c_omp_declare_simd_clauses_to_decls (*decl_p, clauses); - clauses = finish_omp_clauses (clauses); + clauses = finish_omp_clauses (clauses, false); tree parms = DECL_ARGUMENTS (*decl_p); clauses = c_omp_declare_simd_clauses_to_numbers (parms, clauses); @@ -13445,6 +13445,14 @@ tsubst_omp_clauses (tree clauses, bool declare_simd, case OMP_CLAUSE_THREAD_LIMIT: case OMP_CLAUSE_SAFELEN: case OMP_CLAUSE_SIMDLEN: + case OMP_CLAUSE_NUM_GANGS: + case OMP_CLAUSE_NUM_WORKERS: + case OMP_CLAUSE_VECTOR_LENGTH: + case OMP_CLAUSE_GANG: + case OMP_CLAUSE_WORKER: + case OMP_CLAUSE_VECTOR: + case OMP_CLAUSE_ASYNC: + case OMP_CLAUSE_WAIT: OMP_CLAUSE_OPERAND (nc, 0) = tsubst_expr (OMP_CLAUSE_OPERAND (oc, 0), args, complain, in_decl, /*integral_constant_expression_p=*/false); @@ -13491,6 +13499,10 @@ tsubst_omp_clauses (tree clauses, bool declare_simd, case OMP_CLAUSE_PARALLEL: case OMP_CLAUSE_SECTIONS: case OMP_CLAUSE_TASKGROUP: + case OMP_CLAUSE_INDEPENDENT: + case OMP_CLAUSE_AUTO: + case OMP_CLAUSE_SEQ: + case OMP_CLAUSE_TILE: break; default: gcc_unreachable (); @@ -13499,7 +13511,7 @@ tsubst_omp_clauses (tree clauses, bool declare_simd, new_clauses = nreverse (new_clauses); if (!declare_simd) - new_clauses = finish_omp_clauses (new_clauses); + new_clauses = finish_omp_clauses (new_clauses, false); return new_clauses; } @@ -13639,7 +13651,7 @@ tsubst_omp_for_iterator (tree t, int i, tree declv, tree initv, { c = build_omp_clause (input_location, OMP_CLAUSE_PRIVATE); OMP_CLAUSE_DECL (c) = decl; - c = finish_omp_clauses (c); + c = finish_omp_clauses (c, false); if (c) { OMP_CLAUSE_CHAIN (c) = *clauses; @@ -14108,6 +14120,22 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl, = OMP_PARALLEL_COMBINED (t); break; + case OACC_PARALLEL: + tmp = tsubst_omp_clauses (OACC_PARALLEL_CLAUSES (t), false, + args, complain, in_decl); + stmt = begin_omp_parallel (); + RECUR (OACC_PARALLEL_BODY (t)); + finish_oacc_parallel (tmp, stmt); + break; + + case OACC_KERNELS: + tmp = tsubst_omp_clauses (OACC_KERNELS_CLAUSES (t), false, + args, complain, in_decl); + stmt = begin_omp_parallel (); + RECUR (OACC_KERNELS_BODY (t)); + finish_oacc_kernels (tmp, stmt); + break; + case OMP_TASK: tmp = tsubst_omp_clauses (OMP_TASK_CLAUSES (t), false, args, complain, in_decl); @@ -14121,6 +14149,7 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl, case CILK_SIMD: case CILK_FOR: case OMP_DISTRIBUTE: + case OACC_LOOP: { tree clauses, body, pre_body; tree declv = NULL_TREE, initv = NULL_TREE, condv = NULL_TREE; @@ -14186,6 +14215,7 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl, add_stmt (t); break; + case OACC_DATA: case OMP_TARGET_DATA: case OMP_TARGET: tmp = tsubst_omp_clauses (OMP_CLAUSES (t), false, @@ -14203,10 +14233,13 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl, break; case OMP_TARGET_UPDATE: - tmp = tsubst_omp_clauses (OMP_TARGET_UPDATE_CLAUSES (t), false, + case OACC_ENTER_DATA: + case OACC_EXIT_DATA: + case OACC_UPDATE: + tmp = tsubst_omp_clauses (OMP_STANDALONE_CLAUSES (t), false, args, complain, in_decl); t = copy_node (t); - OMP_TARGET_UPDATE_CLAUSES (t) = tmp; + OMP_STANDALONE_CLAUSES (t) = tmp; add_stmt (t); break; diff --git gcc/cp/semantics.c gcc/cp/semantics.c index 0fc08b5f..ada1203 100644 --- gcc/cp/semantics.c +++ gcc/cp/semantics.c @@ -5294,19 +5294,21 @@ finish_omp_reduction_clause (tree c, bool *need_default_ctor, bool *need_dtor) Remove any elements from the list that are invalid. */ tree -finish_omp_clauses (tree clauses) +finish_omp_clauses (tree clauses, bool oacc) { bitmap_head generic_head, firstprivate_head, lastprivate_head; - bitmap_head aligned_head; + bitmap_head aligned_head, oacc_data_head; tree c, t, *pc; bool branch_seen = false; bool copyprivate_seen = false; + bool oacc_data = false; bitmap_obstack_initialize (NULL); bitmap_initialize (&generic_head, &bitmap_default_obstack); bitmap_initialize (&firstprivate_head, &bitmap_default_obstack); bitmap_initialize (&lastprivate_head, &bitmap_default_obstack); bitmap_initialize (&aligned_head, &bitmap_default_obstack); + bitmap_initialize (&oacc_data_head, &bitmap_default_obstack); for (pc = &clauses, c = clauses; c ; c = *pc) { @@ -5317,9 +5319,21 @@ finish_omp_clauses (tree clauses) case OMP_CLAUSE_SHARED: goto check_dup_generic; case OMP_CLAUSE_PRIVATE: - goto check_dup_generic; + if (oacc) + { + oacc_data = true; + goto check_dup_oacc; + } + else + goto check_dup_generic; case OMP_CLAUSE_REDUCTION: - goto check_dup_generic; + if (oacc) + { + oacc_data = false; + goto check_dup_oacc; + } + else + goto check_dup_generic; case OMP_CLAUSE_COPYPRIVATE: copyprivate_seen = true; goto check_dup_generic; @@ -5403,6 +5417,44 @@ finish_omp_clauses (tree clauses) else bitmap_set_bit (&generic_head, DECL_UID (t)); break; + check_dup_oacc: + t = OMP_CLAUSE_DECL (c); + if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL) + { + if (processing_template_decl) + break; + if (DECL_P (t)) + error ("%qD is not a variable in clause %qs", t, + omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + else + error ("%qE is not a variable in clause %qs", t, + omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + remove = true; + } + else if (oacc_data) + { + if (bitmap_bit_p (&oacc_data_head, DECL_UID (t))) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qE appears more than once in data clauses", t); + remove = true; + } + else + bitmap_set_bit (&oacc_data_head, DECL_UID (t)); + } + else + { + if (bitmap_bit_p (&generic_head, DECL_UID (t))) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qE appears more than once in data clauses", t); + remove = true; + } + else + bitmap_set_bit (&generic_head, DECL_UID (t)); + } + break; + case OMP_CLAUSE_FIRSTPRIVATE: t = OMP_CLAUSE_DECL (c); @@ -5426,6 +5478,37 @@ finish_omp_clauses (tree clauses) bitmap_set_bit (&firstprivate_head, DECL_UID (t)); break; + case OMP_CLAUSE_GANG: + case OMP_CLAUSE_VECTOR: + case OMP_CLAUSE_WORKER: + /* Operand 0 is the num: or length: argument. */ + t = OMP_CLAUSE_OPERAND (c, 0); + if (t == NULL_TREE) + break; + + t = maybe_convert_cond (t); + if (t == error_mark_node) + remove = true; + else if (!processing_template_decl) + t = fold_build_cleanup_point_expr (TREE_TYPE (t), t); + OMP_CLAUSE_OPERAND (c, 0) = t; + + if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_GANG) + break; + + /* Ooperand 1 is the gang static: argument. */ + t = OMP_CLAUSE_OPERAND (c, 1); + if (t == NULL_TREE) + break; + + t = maybe_convert_cond (t); + if (t == error_mark_node) + remove = true; + else if (!processing_template_decl) + t = fold_build_cleanup_point_expr (TREE_TYPE (t), t); + OMP_CLAUSE_OPERAND (c, 1) = t; + break; + case OMP_CLAUSE_LASTPRIVATE: t = OMP_CLAUSE_DECL (c); if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL) @@ -5469,13 +5552,30 @@ finish_omp_clauses (tree clauses) break; case OMP_CLAUSE_NUM_THREADS: - t = OMP_CLAUSE_NUM_THREADS_EXPR (c); + case OMP_CLAUSE_NUM_GANGS: + case OMP_CLAUSE_NUM_WORKERS: + case OMP_CLAUSE_VECTOR_LENGTH: + t = OMP_CLAUSE_OPERAND (c, 0); if (t == error_mark_node) remove = true; else if (!type_dependent_expression_p (t) && !INTEGRAL_TYPE_P (TREE_TYPE (t))) { - error ("num_threads expression must be integral"); + switch (OMP_CLAUSE_CODE (c)) + { + case OMP_CLAUSE_NUM_THREADS: + error ("num_threads expression must be integral"); break; + case OMP_CLAUSE_NUM_GANGS: + error ("% expression must be integral"); break; + case OMP_CLAUSE_NUM_WORKERS: + error ("% expression must be integral"); + break; + case OMP_CLAUSE_VECTOR_LENGTH: + error ("% expression must be integral"); + break; + default: + error ("invalid argument"); + } remove = true; } else @@ -5483,7 +5583,7 @@ finish_omp_clauses (tree clauses) t = mark_rvalue_use (t); if (!processing_template_decl) t = fold_build_cleanup_point_expr (TREE_TYPE (t), t); - OMP_CLAUSE_NUM_THREADS_EXPR (c) = t; + OMP_CLAUSE_OPERAND (c, 0) = t; } break; @@ -5591,16 +5691,6 @@ finish_omp_clauses (tree clauses) } break; - case OMP_CLAUSE_VECTOR_LENGTH: - t = OMP_CLAUSE_VECTOR_LENGTH_EXPR (c); - t = maybe_convert_cond (t); - if (t == error_mark_node) - remove = true; - else if (!processing_template_decl) - t = fold_build_cleanup_point_expr (TREE_TYPE (t), t); - OMP_CLAUSE_VECTOR_LENGTH_EXPR (c) = t; - break; - case OMP_CLAUSE_WAIT: t = OMP_CLAUSE_WAIT_EXPR (c); if (t == error_mark_node) @@ -5861,6 +5951,13 @@ finish_omp_clauses (tree clauses) case OMP_CLAUSE_TASKGROUP: case OMP_CLAUSE_PROC_BIND: case OMP_CLAUSE__CILK_FOR_COUNT_: + case OMP_CLAUSE_USE_DEVICE: + case OMP_CLAUSE_AUTO: + case OMP_CLAUSE_INDEPENDENT: + case OMP_CLAUSE_SEQ: + case OMP_CLAUSE_BIND: + case OMP_CLAUSE_NOHOST: + case OMP_CLAUSE_TILE: break; case OMP_CLAUSE_INBRANCH: @@ -6146,6 +6243,24 @@ finish_oacc_data (tree clauses, tree block) return add_stmt (stmt); } +/* Generate OACC_HOST_DATA, with CLAUSES and BLOCK as its compound + statement. */ + +tree +finish_oacc_host_data (tree clauses, tree block) +{ + tree stmt; + + block = finish_omp_structured_block (block); + + stmt = make_node (OACC_HOST_DATA); + TREE_TYPE (stmt) = void_type_node; + OACC_HOST_DATA_CLAUSES (stmt) = clauses; + OACC_HOST_DATA_BODY (stmt) = block; + + return add_stmt (stmt); +} + /* Generate OACC_KERNELS, with CLAUSES and BLOCK as its compound statement. LOC is the location of the OACC_KERNELS. */ @@ -6805,7 +6920,7 @@ finish_omp_for (location_t locus, enum tree_code code, tree declv, tree initv, OMP_CLAUSE_OPERAND (c, 0) = cilk_for_number_of_iterations (omp_for); OMP_CLAUSE_CHAIN (c) = clauses; - OMP_PARALLEL_CLAUSES (omp_par) = finish_omp_clauses (c); + OMP_PARALLEL_CLAUSES (omp_par) = finish_omp_clauses (c, false); add_stmt (omp_par); return omp_par; } Grüße, Thomas