2016-11-22 Cesar Philippidis <cesar@codesourcery.com>
Thomas Schwinge <thomas@codesourcery.com>
gcc/cp/
* cp-tree.h (bind_decls_match): Declare.
* decl.c (bind_decls_match): New function.
* parser.c (cp_parser_omp_clause_name):
(cp_parser_oacc_simple_clause): Remove unused cp_parser argument.
(cp_parser_oacc_shape_clause): New location_t loc argument. Use it
to report more accurate diagnostics. Remove parser argument.
(cp_parser_oacc_clause_bind): New function.
(cp_parser_oacc_all_clauses): Handle OpenACC bind and nohost clauses.
Update calls to c_parser_oacc_{simple,shape}_clause.
(OACC_ROUTINE_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_{BIND,NOHOST}.
(cp_parser_oacc_routine): Update diagnostics.
(cp_parser_late_parsing_oacc_routine): Likewise.
(cp_finalize_oacc_routine): Likewise.
* semantics.c (finish_omp_clauses): Handle OMP_CLAUSE_{BIND,NOHOST}.
@@ -5785,6 +5785,7 @@ extern void finish_scope (void);
extern void push_switch (tree);
extern void pop_switch (void);
extern tree make_lambda_name (void);
+extern int bind_decls_match (tree, tree);
extern int decls_match (tree, tree);
extern tree duplicate_decls (tree, tree, bool);
extern tree declare_local_label (tree);
@@ -1198,6 +1198,138 @@ decls_match (tree newdecl, tree olddecl)
return types_match;
}
+/* Similiar to decls_match, but only applies to FUNCTION_DECLS. Functions
+ in separate namespaces may match.
+*/
+
+int
+bind_decls_match (tree newdecl, tree olddecl)
+{
+ int types_match;
+
+ if (newdecl == olddecl)
+ return 1;
+
+ if (TREE_CODE (newdecl) != TREE_CODE (olddecl))
+ /* If the two DECLs are not even the same kind of thing, we're not
+ interested in their types. */
+ return 0;
+
+ gcc_assert (DECL_P (newdecl));
+ gcc_assert (TREE_CODE (newdecl) == FUNCTION_DECL);
+
+ tree f1 = TREE_TYPE (newdecl);
+ tree f2 = TREE_TYPE (olddecl);
+ tree p1 = TYPE_ARG_TYPES (f1);
+ tree p2 = TYPE_ARG_TYPES (f2);
+ tree r2;
+
+ /* Specializations of different templates are different functions
+ even if they have the same type. */
+ tree t1 = (DECL_USE_TEMPLATE (newdecl)
+ ? DECL_TI_TEMPLATE (newdecl)
+ : NULL_TREE);
+ tree t2 = (DECL_USE_TEMPLATE (olddecl)
+ ? DECL_TI_TEMPLATE (olddecl)
+ : NULL_TREE);
+ if (t1 != t2)
+ return 0;
+
+ if (CP_DECL_CONTEXT (newdecl) != CP_DECL_CONTEXT (olddecl)
+ && TREE_CODE (CP_DECL_CONTEXT (newdecl)) != NAMESPACE_DECL
+ && TREE_CODE (CP_DECL_CONTEXT (olddecl)) != NAMESPACE_DECL
+ && ! (DECL_EXTERN_C_P (newdecl)
+ && DECL_EXTERN_C_P (olddecl)))
+ return 0;
+
+ /* A new declaration doesn't match a built-in one unless it
+ is also extern "C". */
+ if (DECL_IS_BUILTIN (olddecl)
+ && DECL_EXTERN_C_P (olddecl) && !DECL_EXTERN_C_P (newdecl))
+ return 0;
+
+ if (TREE_CODE (f1) != TREE_CODE (f2))
+ return 0;
+
+ /* A declaration with deduced return type should use its pre-deduction
+ type for declaration matching. */
+ r2 = fndecl_declared_return_type (olddecl);
+
+ if (same_type_p (TREE_TYPE (f1), r2))
+ {
+ if (!prototype_p (f2) && DECL_EXTERN_C_P (olddecl)
+ && (DECL_BUILT_IN (olddecl)
+#ifndef NO_IMPLICIT_EXTERN_C
+ || (DECL_IN_SYSTEM_HEADER (newdecl) && !DECL_CLASS_SCOPE_P (newdecl))
+ || (DECL_IN_SYSTEM_HEADER (olddecl) && !DECL_CLASS_SCOPE_P (olddecl))
+#endif
+ ))
+ {
+ types_match = self_promoting_args_p (p1);
+ if (p1 == void_list_node)
+ TREE_TYPE (newdecl) = TREE_TYPE (olddecl);
+ }
+#ifndef NO_IMPLICIT_EXTERN_C
+ else if (!prototype_p (f1)
+ && (DECL_EXTERN_C_P (olddecl)
+ && DECL_IN_SYSTEM_HEADER (olddecl)
+ && !DECL_CLASS_SCOPE_P (olddecl))
+ && (DECL_EXTERN_C_P (newdecl)
+ && DECL_IN_SYSTEM_HEADER (newdecl)
+ && !DECL_CLASS_SCOPE_P (newdecl)))
+ {
+ types_match = self_promoting_args_p (p2);
+ TREE_TYPE (newdecl) = TREE_TYPE (olddecl);
+ }
+#endif
+ else
+ types_match =
+ compparms (p1, p2)
+ && type_memfn_rqual (f1) == type_memfn_rqual (f2)
+ && (TYPE_ATTRIBUTES (TREE_TYPE (newdecl)) == NULL_TREE
+ || comp_type_attributes (TREE_TYPE (newdecl),
+ TREE_TYPE (olddecl)) != 0);
+ }
+ else
+ types_match = 0;
+
+ /* The decls dont match if they correspond to two different versions
+ of the same function. Disallow extern "C" functions to be
+ versions for now. */
+ if (types_match
+ && !DECL_EXTERN_C_P (newdecl)
+ && !DECL_EXTERN_C_P (olddecl)
+ && targetm.target_option.function_versions (newdecl, olddecl))
+ {
+ /* Mark functions as versions if necessary. Modify the mangled decl
+ name if necessary. */
+ if (DECL_FUNCTION_VERSIONED (newdecl)
+ && DECL_FUNCTION_VERSIONED (olddecl))
+ return 0;
+ if (!DECL_FUNCTION_VERSIONED (newdecl))
+ {
+ DECL_FUNCTION_VERSIONED (newdecl) = 1;
+ if (DECL_ASSEMBLER_NAME_SET_P (newdecl))
+ mangle_decl (newdecl);
+ }
+ if (!DECL_FUNCTION_VERSIONED (olddecl))
+ {
+ DECL_FUNCTION_VERSIONED (olddecl) = 1;
+ if (DECL_ASSEMBLER_NAME_SET_P (olddecl))
+ mangle_decl (olddecl);
+ }
+ cgraph_node::record_function_versions (olddecl, newdecl);
+ return 0;
+ }
+
+ /* Normal functions can be constrained, as can variable partial
+ specializations. */
+ if (types_match && VAR_OR_FUNCTION_DECL_P (newdecl))
+ types_match = equivalently_constrained (newdecl, olddecl);
+
+ return types_match;
+}
+
/* If NEWDECL is `static' and an `extern' was seen previously,
warn about it. OLDDECL is the previous declaration.
@@ -30497,6 +30497,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;
@@ -30576,6 +30580,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))
@@ -30991,8 +30997,7 @@ cp_parser_oacc_data_clause_deviceptr (cp_parser *parser, tree list)
seq */
static tree
-cp_parser_oacc_simple_clause (cp_parser * /* parser */,
- enum omp_clause_code code,
+cp_parser_oacc_simple_clause (enum omp_clause_code code,
tree list, location_t location)
{
check_no_duplicate_clause (list, code, omp_clause_code_name[code], location);
@@ -31052,13 +31057,13 @@ cp_parser_oacc_single_int_clause (cp_parser *parser, omp_clause_code code,
*/
static tree
-cp_parser_oacc_shape_clause (cp_parser *parser, omp_clause_code kind,
+cp_parser_oacc_shape_clause (cp_parser *parser, location_t loc,
+ omp_clause_code kind,
const char *str, tree list)
{
const char *id = "num";
cp_lexer *lexer = parser->lexer;
tree ops[2] = { NULL_TREE, NULL_TREE }, c;
- location_t loc = cp_lexer_peek_token (lexer)->location;
if (kind == OMP_CLAUSE_VECTOR)
id = "length";
@@ -32845,6 +32850,82 @@ 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;
+
+ check_no_duplicate_clause (list, OMP_CLAUSE_BIND, "bind", loc);
+
+ 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;
+ }
+ tree name = error_mark_node;
+ cp_token *token = cp_lexer_peek_token (parser->lexer);
+ if (cp_lexer_next_token_is (parser->lexer, CPP_NAME))
+ {
+ tree id = cp_parser_id_expression (parser, /*template_p=*/false,
+ /*check_dependency_p=*/true,
+ /*template_p=*/NULL,
+ /*declarator_p=*/false,
+ /*optional_p=*/false);
+ tree decl = cp_parser_lookup_name_simple (parser, id, token->location);
+ if (id != error_mark_node && decl == error_mark_node)
+ cp_parser_name_lookup_error (parser, id, decl, NLE_NULL,
+ token->location);
+ if (!decl || decl == error_mark_node)
+ error_at (token->location, "%qE has not been declared",
+ token->u.value);
+ else if (is_overloaded_fn (decl)
+ && (TREE_CODE (decl) != FUNCTION_DECL
+ || DECL_FUNCTION_TEMPLATE_P (decl)))
+ error_at (token->location, "%qE names a set of overloads",
+ token->u.value);
+ else if (TREE_CODE (decl) != FUNCTION_DECL)
+ error_at (token->location,
+ "%qE does not refer to a function",
+ token->u.value);
+ else
+ name = decl;
+ }
+ else if (cp_lexer_next_token_is (parser->lexer, CPP_STRING))
+ {
+ name = token->u.value;
+ cp_lexer_consume_token (parser->lexer);
+
+ /* This shouldn't be an empty string. */
+ if (strcmp (TREE_STRING_POINTER (name), "\"\"") == 0)
+ error_at (token->location,
+ "bind argument must not be an empty string");
+
+ parser->translate_strings_p = save_translate_strings_p;
+ }
+ else
+ {
+ cp_parser_error (parser,
+ "expected identifier or character string literal");
+ cp_parser_skip_to_closing_parenthesis (parser, false, false, true);
+ return list;
+ }
+ cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN);
+ if (name != error_mark_node)
+ {
+ tree c = build_omp_clause (loc, OMP_CLAUSE_BIND);
+ OMP_CLAUSE_BIND_NAME (c) = name;
+ OMP_CLAUSE_CHAIN (c) = list;
+ list = c;
+ }
+ return list;
+}
+
/* Parse all OpenACC clauses. The set clauses allowed by the directive
is a bitmask in MASK. Return the list of clauses found. */
@@ -32877,10 +32958,14 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
c_name = "async";
break;
case PRAGMA_OACC_CLAUSE_AUTO:
- clauses = cp_parser_oacc_simple_clause (parser, OMP_CLAUSE_AUTO,
- clauses, here);
+ clauses = cp_parser_oacc_simple_clause (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";
@@ -32928,7 +33013,7 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
break;
case PRAGMA_OACC_CLAUSE_GANG:
c_name = "gang";
- clauses = cp_parser_oacc_shape_clause (parser, OMP_CLAUSE_GANG,
+ clauses = cp_parser_oacc_shape_clause (parser, here, OMP_CLAUSE_GANG,
c_name, clauses);
break;
case PRAGMA_OACC_CLAUSE_HOST:
@@ -32940,8 +33025,7 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
c_name = "if";
break;
case PRAGMA_OACC_CLAUSE_INDEPENDENT:
- clauses = cp_parser_oacc_simple_clause (parser,
- OMP_CLAUSE_INDEPENDENT,
+ clauses = cp_parser_oacc_simple_clause (OMP_CLAUSE_INDEPENDENT,
clauses, here);
c_name = "independent";
break;
@@ -32949,6 +33033,11 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "link";
break;
+ case PRAGMA_OACC_CLAUSE_NOHOST:
+ clauses = cp_parser_oacc_simple_clause (OMP_CLAUSE_NOHOST, clauses,
+ here);
+ c_name = "nohost";
+ break;
case PRAGMA_OACC_CLAUSE_NUM_GANGS:
code = OMP_CLAUSE_NUM_GANGS;
c_name = "num_gangs";
@@ -32995,8 +33084,8 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
c_name = "self";
break;
case PRAGMA_OACC_CLAUSE_SEQ:
- clauses = cp_parser_oacc_simple_clause (parser, OMP_CLAUSE_SEQ,
- clauses, here);
+ clauses = cp_parser_oacc_simple_clause (OMP_CLAUSE_SEQ, clauses,
+ here);
c_name = "seq";
break;
case PRAGMA_OACC_CLAUSE_TILE:
@@ -33010,7 +33099,8 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
break;
case PRAGMA_OACC_CLAUSE_VECTOR:
c_name = "vector";
- clauses = cp_parser_oacc_shape_clause (parser, OMP_CLAUSE_VECTOR,
+ clauses = cp_parser_oacc_shape_clause (parser, here,
+ OMP_CLAUSE_VECTOR,
c_name, clauses);
break;
case PRAGMA_OACC_CLAUSE_VECTOR_LENGTH:
@@ -33025,7 +33115,8 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
break;
case PRAGMA_OACC_CLAUSE_WORKER:
c_name = "worker";
- clauses = cp_parser_oacc_shape_clause (parser, OMP_CLAUSE_WORKER,
+ clauses = cp_parser_oacc_shape_clause (parser, here,
+ OMP_CLAUSE_WORKER,
c_name, clauses);
break;
default:
@@ -37266,7 +37357,10 @@ cp_parser_omp_taskloop (cp_parser *parser, cp_token *pragma_tok,
( (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_SEQ) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_BIND) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NOHOST))
+
/* Parse the OpenACC routine pragma. This has an optional '( name )'
@@ -37325,6 +37419,9 @@ cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok,
= cp_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK,
"#pragma acc routine",
cp_lexer_peek_token (parser->lexer));
+ /* The clauses are in reverse order; fix that to make later diagnostic
+ emission easier. */
+ data.clauses = nreverse (data.clauses);
if (decl && is_overloaded_fn (decl)
&& (TREE_CODE (decl) != FUNCTION_DECL
@@ -37336,16 +37433,6 @@ cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok,
return;
}
- /* Perhaps we should use the same rule as declarations in different
- namespaces? */
- if (!DECL_NAMESPACE_SCOPE_P (decl))
- {
- error_at (name_loc,
- "%qD does not refer to a namespace scope function", decl);
- parser->oacc_routine = NULL;
- return;
- }
-
if (TREE_CODE (decl) != FUNCTION_DECL)
{
error_at (name_loc, "%qD does not refer to a function", decl);
@@ -37421,6 +37508,9 @@ cp_parser_late_parsing_oacc_routine (cp_parser *parser, tree attrs)
parser->oacc_routine->clauses
= cp_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK,
"#pragma acc routine", pragma_tok);
+ /* The clauses are in reverse order; fix that to make later diagnostic
+ emission easier. */
+ parser->oacc_routine->clauses = nreverse (parser->oacc_routine->clauses);
cp_parser_pop_lexer (parser);
/* Later, cp_finalize_oacc_routine will process the clauses, and then set
fndecl_seen. */
@@ -37455,31 +37545,70 @@ cp_finalize_oacc_routine (cp_parser *parser, tree fndecl, bool is_defn)
return;
}
- if (get_oacc_fn_attrib (fndecl))
+ /* Process the bind clause, if present. */
+ for (tree c = parser->oacc_routine->clauses;
+ c;
+ c = OMP_CLAUSE_CHAIN (c))
{
- error_at (parser->oacc_routine->loc,
- "%<#pragma acc routine%> already applied to %qD", fndecl);
- parser->oacc_routine = NULL;
- return;
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_BIND)
+ continue;
+
+ tree bind_decl = OMP_CLAUSE_BIND_NAME (c);
+
+ /* String arguments don't require any special treatment. */
+ if (TREE_CODE (bind_decl) != FUNCTION_DECL)
+ break;
+
+ if (!bind_decls_match (bind_decl, fndecl))
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "bind identifier %qE is not compatible with "
+ "function %qE", bind_decl, fndecl);
+ parser->oacc_routine = NULL;
+ return;
+ }
+
+ tree name_id = decl_assembler_name (bind_decl);
+ tree name = build_string (IDENTIFIER_LENGTH (name_id),
+ IDENTIFIER_POINTER (name_id));
+ OMP_CLAUSE_BIND_NAME (c) = name;
+
+ break;
}
- if (TREE_USED (fndecl) || (!is_defn && DECL_SAVED_TREE (fndecl)))
+ int compatible
+ = verify_oacc_routine_clauses (fndecl, &parser->oacc_routine->clauses,
+ parser->oacc_routine->loc,
+ "#pragma acc routine");
+ if (compatible < 0)
{
- error_at (parser->oacc_routine->loc,
- "%<#pragma acc routine%> must be applied before %s",
- TREE_USED (fndecl) ? "use" : "definition");
parser->oacc_routine = NULL;
return;
}
+ if (compatible > 0)
+ {
+ }
+ else
+ {
+ if (TREE_USED (fndecl) || (!is_defn && DECL_SAVED_TREE (fndecl)))
+ {
+ error_at (parser->oacc_routine->loc,
+ "%<#pragma acc routine%> must be applied before %s",
+ TREE_USED (fndecl) ? "use" : "definition");
+ parser->oacc_routine = NULL;
+ return;
+ }
- /* Process the routine's dimension clauses. */
- tree dims = build_oacc_routine_dims (parser->oacc_routine->clauses);
- replace_oacc_fn_attrib (fndecl, dims);
-
- /* Add an "omp declare target" attribute. */
- DECL_ATTRIBUTES (fndecl)
- = tree_cons (get_identifier ("omp declare target"),
- NULL_TREE, DECL_ATTRIBUTES (fndecl));
+ /* Set the routine's level of parallelism. */
+ tree dims = build_oacc_routine_dims (parser->oacc_routine->clauses);
+ replace_oacc_fn_attrib (fndecl, dims);
+
+ /* Add an "omp declare target" attribute. */
+ DECL_ATTRIBUTES (fndecl)
+ = tree_cons (get_identifier ("omp declare target"),
+ parser->oacc_routine->clauses,
+ DECL_ATTRIBUTES (fndecl));
+ }
/* Don't unset parser->oacc_routine here: we may still need it to
diagnose wrong usage. But, remember that we've used this "#pragma acc
@@ -7074,6 +7074,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
case OMP_CLAUSE_AUTO:
case OMP_CLAUSE_INDEPENDENT:
case OMP_CLAUSE_SEQ:
+ case OMP_CLAUSE_BIND:
+ case OMP_CLAUSE_NOHOST:
break;
case OMP_CLAUSE_TILE: