2016-11-11 Cesar Philippidis <cesar@codesourcery.com>
Thomas Schwinge <thomas@codesourcery.com>
gcc/c-family/
* c-attribs.c (c_common_attribute_table): Adjust "omp declare target".
* c-pragma.h (enum pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_BIND
and PRAGMA_OACC_CLAUSE_NOHOST.
gcc/
* gimplify.c (gimplify_scan_omp_clauses): Handle OMP_CLAUSE_BIND and
OMP_CLAUSE_NOHOST.
(gimplify_adjust_omp_clauses): Likewise.
* omp-low.c (scan_sharing_clauses): Likewise.
(verify_oacc_routine_clauses): New function.
(maybe_discard_oacc_function): New function.
(execute_oacc_device_lower): Don't generate code for NOHOST.
* omp-low.h (verify_oacc_routine_clauses): Declare.
* tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_BIND and
OMP_CLAUSE_NOHOST.
* tree-pretty-print.c (dump_omp_clause): Likewise.
* tree.c (omp_clause_num_ops): Likewise.
(omp_clause_code_name): Likewise.
(walk_tree_1): Handle OMP_CLAUSE_BIND, OMP_CLAUSE_NOHOST.
* tree.h (OMP_CLAUSE_BIND_NAME): Define.
@@ -322,7 +322,7 @@ const struct attribute_spec c_common_attribute_table[] =
handle_omp_declare_simd_attribute, false },
{ "simd", 0, 1, true, false, false,
handle_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 },
{ "omp declare target link", 0, 0, true, false, false,
handle_omp_declare_target_attribute, false },
@@ -149,6 +149,7 @@ 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,
@@ -158,6 +159,7 @@ enum pragma_omp_clause {
PRAGMA_OACC_CLAUSE_GANG,
PRAGMA_OACC_CLAUSE_HOST,
PRAGMA_OACC_CLAUSE_INDEPENDENT,
+ PRAGMA_OACC_CLAUSE_NOHOST,
PRAGMA_OACC_CLAUSE_NUM_GANGS,
PRAGMA_OACC_CLAUSE_NUM_WORKERS,
PRAGMA_OACC_CLAUSE_PRESENT,
@@ -8373,6 +8373,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
ctx->default_kind = OMP_CLAUSE_DEFAULT_KIND (c);
break;
+ case OMP_CLAUSE_BIND:
+ case OMP_CLAUSE_NOHOST:
default:
gcc_unreachable ();
}
@@ -9112,6 +9114,8 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
remove = true;
break;
+ case OMP_CLAUSE_BIND:
+ case OMP_CLAUSE_NOHOST:
default:
gcc_unreachable ();
}
@@ -2201,6 +2201,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
install_var_local (decl, ctx);
break;
+ case OMP_CLAUSE_BIND:
+ case OMP_CLAUSE_NOHOST:
case OMP_CLAUSE_TILE:
case OMP_CLAUSE__CACHE_:
default:
@@ -2365,6 +2367,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
case OMP_CLAUSE__GRIDDIM_:
break;
+ case OMP_CLAUSE_BIND:
+ case OMP_CLAUSE_NOHOST:
case OMP_CLAUSE_TILE:
case OMP_CLAUSE__CACHE_:
default:
@@ -12684,9 +12688,192 @@ set_oacc_fn_attrib (tree fn, tree clauses, bool is_kernel, vec<tree> *args)
}
}
-/* Process the routine's dimension clauess to generate an attribute
- value. Issue diagnostics as appropriate. We default to SEQ
- (OpenACC 2.5 clarifies this). All dimensions have a size of zero
+/* Verify OpenACC routine clauses.
+
+ Returns 0 if FNDECL should be marked as an accelerator routine, 1 if it has
+ already been marked in compatible way, and -1 if incompatible. Upon
+ returning, the chain of clauses will contain exactly one clause specifying
+ the level of parallelism. */
+
+int
+verify_oacc_routine_clauses (tree fndecl, tree *clauses, location_t loc,
+ const char *routine_str)
+{
+ tree c_level = NULL_TREE;
+ tree c_bind = NULL_TREE;
+ tree c_nohost = NULL_TREE;
+ tree c_p = NULL_TREE;
+ for (tree c = *clauses; c; c_p = 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:
+ if (c_level == NULL_TREE)
+ c_level = c;
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_CODE (c_level))
+ {
+ /* This has already been diagnosed in the front ends. */
+ /* Drop the duplicate clause. */
+ gcc_checking_assert (c_p != NULL_TREE);
+ OMP_CLAUSE_CHAIN (c_p) = OMP_CLAUSE_CHAIN (c);
+ c = c_p;
+ }
+ else
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%qs specifies a conflicting level of parallelism",
+ omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+ inform (OMP_CLAUSE_LOCATION (c_level),
+ "... to the previous %qs clause here",
+ omp_clause_code_name[OMP_CLAUSE_CODE (c_level)]);
+ /* Drop the conflicting clause. */
+ gcc_checking_assert (c_p != NULL_TREE);
+ OMP_CLAUSE_CHAIN (c_p) = OMP_CLAUSE_CHAIN (c);
+ c = c_p;
+ }
+ break;
+ case OMP_CLAUSE_BIND:
+ /* Don't bother with duplicate clauses at this point. */
+ c_bind = c;
+ break;
+ case OMP_CLAUSE_NOHOST:
+ /* Don't bother with duplicate clauses at this point. */
+ c_nohost = c;
+ break;
+ default:
+ gcc_unreachable ();
+ }
+ if (c_level == NULL_TREE)
+ {
+ /* OpenACC 2.5 makes this an error; for the current OpenACC 2.0a
+ implementation add an implicit "seq" clause. */
+ c_level = build_omp_clause (loc, OMP_CLAUSE_SEQ);
+ OMP_CLAUSE_CHAIN (c_level) = *clauses;
+ *clauses = c_level;
+ }
+ /* In *clauses, we now have exactly one clause specifying the level of
+ parallelism. */
+
+ /* Still got some work to do for Fortran... */
+ if (fndecl == NULL_TREE)
+ return 0;
+
+ tree attr
+ = lookup_attribute ("omp declare target", DECL_ATTRIBUTES (fndecl));
+ if (attr != NULL_TREE)
+ {
+ /* If a "#pragma acc routine" has already been applied, just verify
+ this one for compatibility. */
+ /* Collect previous directive's clauses. */
+ tree c_level_p = NULL_TREE;
+ tree c_bind_p = NULL_TREE;
+ tree c_nohost_p = NULL_TREE;
+ for (tree c = TREE_VALUE (attr); 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:
+ gcc_checking_assert (c_level_p == NULL_TREE);
+ c_level_p = c;
+ break;
+ case OMP_CLAUSE_BIND:
+ /* Don't bother with duplicate clauses at this point. */
+ c_bind_p = c;
+ break;
+ case OMP_CLAUSE_NOHOST:
+ /* Don't bother with duplicate clauses at this point. */
+ c_nohost_p = c;
+ break;
+ default:
+ gcc_unreachable ();
+ }
+ gcc_checking_assert (c_level_p != NULL_TREE);
+ /* ..., and compare to current directive's, which we've already collected
+ above. */
+ tree c_diag;
+ tree c_diag_p;
+ /* Matching level of parallelism? */
+ if (OMP_CLAUSE_CODE (c_level) != OMP_CLAUSE_CODE (c_level_p))
+ {
+ c_diag = c_level;
+ c_diag_p = c_level_p;
+ goto incompatible;
+ }
+ /* Matching bind clauses? */
+ if ((c_bind == NULL_TREE) != (c_bind_p == NULL_TREE))
+ {
+ c_diag = c_bind;
+ c_diag_p = c_bind_p;
+ goto incompatible;
+ }
+ /* Matching bind clauses' names? */
+ if ((c_bind != NULL_TREE) && (c_bind_p != NULL_TREE))
+ {
+ tree c_bind_name = OMP_CLAUSE_BIND_NAME (c_bind);
+ tree c_bind_name_p = OMP_CLAUSE_BIND_NAME (c_bind_p);
+ /* TODO: will/should actually be the trees/strings/string pointers be
+ identical? */
+ if (strcmp (TREE_STRING_POINTER (c_bind_name),
+ TREE_STRING_POINTER (c_bind_name_p)) != 0)
+ {
+ c_diag = c_bind;
+ c_diag_p = c_bind_p;
+ goto incompatible;
+ }
+ }
+ /* Matching nohost clauses? */
+ if ((c_nohost == NULL_TREE) != (c_nohost_p == NULL_TREE))
+ {
+ c_diag = c_nohost;
+ c_diag_p = c_nohost_p;
+ goto incompatible;
+ }
+ /* Compatible. */
+ return 1;
+
+ incompatible:
+ if (c_diag != NULL_TREE)
+ error_at (OMP_CLAUSE_LOCATION (c_diag),
+ "incompatible %qs clause when applying"
+ " %<%s%> to %qD, which has already been"
+ " marked as an accelerator routine",
+ omp_clause_code_name[OMP_CLAUSE_CODE (c_diag)],
+ routine_str, fndecl);
+ else if (c_diag_p != NULL_TREE)
+ error_at (loc,
+ "missing %qs clause when applying"
+ " %<%s%> to %qD, which has already been"
+ " marked as an accelerator routine",
+ omp_clause_code_name[OMP_CLAUSE_CODE (c_diag_p)],
+ routine_str, fndecl);
+ else
+ gcc_unreachable ();
+ if (c_diag_p != NULL_TREE)
+ inform (OMP_CLAUSE_LOCATION (c_diag_p),
+ "... with %qs clause here",
+ omp_clause_code_name[OMP_CLAUSE_CODE (c_diag_p)]);
+ else
+ {
+ /* In the front ends, we don't preserve location information for the
+ OpenACC routine directive itself. However, that of c_level_p
+ should be close. */
+ location_t loc_routine = OMP_CLAUSE_LOCATION (c_level_p);
+ inform (loc_routine, "... without %qs clause near to here",
+ omp_clause_code_name[OMP_CLAUSE_CODE (c_diag)]);
+ }
+ /* Incompatible. */
+ return -1;
+ }
+
+ return 0;
+}
+
+/* Process the OpenACC routine's clauses to generate an attribute
+ for the level of parallelism. All dimensions have a size of zero
(dynamic). TREE_PURPOSE is set to indicate whether that dimension
can have a loop partitioned on it. non-zero indicates
yes, zero indicates no. By construction once a non-zero has been
@@ -19694,6 +19881,28 @@ default_goacc_reduction (gcall *call)
gsi_replace_with_seq (&gsi, seq, true);
}
+/* Determine whether DECL should be discarded in this offload
+ compilation. */
+
+static bool
+maybe_discard_oacc_function (tree decl)
+{
+ tree attr = lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl));
+
+ if (!attr)
+ return false;
+
+ enum omp_clause_code kind = OMP_CLAUSE_NOHOST;
+
+#ifdef ACCEL_COMPILER
+ kind = OMP_CLAUSE_BIND;
+#endif
+ if (find_omp_clause (TREE_VALUE (attr), kind))
+ return true;
+
+ return false;
+}
+
/* Main entry point for oacc transformations which run on the device
compiler after LTO, so we know what the target device is at this
point (including the host fallback). */
@@ -19707,6 +19916,14 @@ execute_oacc_device_lower ()
/* Not an offloaded function. */
return 0;
+ if (maybe_discard_oacc_function (current_function_decl))
+ {
+ if (dump_file)
+ fprintf (dump_file, "Discarding function\n");
+ TREE_ASM_WRITTEN (current_function_decl) = 1;
+ return TODO_discard_function;
+ }
+
/* Parse the default dim argument exactly once. */
if ((const void *)flag_openacc_dims != &flag_openacc_dims)
{
@@ -31,6 +31,7 @@ extern bool make_gimple_omp_edges (basic_block, struct omp_region **, int *);
extern void omp_finish_file (void);
extern tree omp_member_access_dummy_var (tree);
extern void replace_oacc_fn_attrib (tree, tree);
+extern int verify_oacc_routine_clauses (tree, tree *, location_t, const char *);
extern tree build_oacc_routine_dims (tree);
extern tree get_oacc_fn_attrib (tree);
extern void set_oacc_fn_attrib (tree, tree, bool, vec<tree> *);
@@ -465,7 +465,13 @@ enum omp_clause_code {
/* OpenMP internal-only clause to specify grid dimensions of a gridified
kernel. */
- OMP_CLAUSE__GRIDDIM_
+ OMP_CLAUSE__GRIDDIM_,
+
+ /* OpenACC clause: bind (string). */
+ OMP_CLAUSE_BIND,
+
+ /* OpenACC clause: nohost. */
+ OMP_CLAUSE_NOHOST
};
#undef DEFTREESTRUCT
@@ -1021,6 +1021,15 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, int flags)
spc, flags, false);
pp_right_paren (pp);
break;
+ case OMP_CLAUSE_NOHOST:
+ pp_string (pp, "nohost");
+ break;
+ case OMP_CLAUSE_BIND:
+ pp_string (pp, "bind(");
+ dump_generic_node (pp, OMP_CLAUSE_BIND_NAME (clause),
+ spc, flags, false);
+ pp_string (pp, ")");
+ break;
case OMP_CLAUSE__GRIDDIM_:
pp_string (pp, "_griddim_(");
@@ -329,6 +329,8 @@ unsigned const char omp_clause_num_ops[] =
1, /* OMP_CLAUSE_VECTOR_LENGTH */
1, /* OMP_CLAUSE_TILE */
2, /* OMP_CLAUSE__GRIDDIM_ */
+ 1, /* OMP_CLAUSE_BIND */
+ 0, /* OMP_CLAUSE_NOHOST */
};
const char * const omp_clause_code_name[] =
@@ -399,7 +401,9 @@ const char * const omp_clause_code_name[] =
"num_workers",
"vector_length",
"tile",
- "_griddim_"
+ "_griddim_",
+ "bind",
+ "nohost",
};
@@ -11869,6 +11873,7 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,
case OMP_CLAUSE__LOOPTEMP_:
case OMP_CLAUSE__SIMDUID_:
case OMP_CLAUSE__CILK_FOR_COUNT_:
+ case OMP_CLAUSE_BIND:
WALK_SUBTREE (OMP_CLAUSE_OPERAND (*tp, 0));
/* FALLTHRU */
@@ -11890,6 +11895,7 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,
case OMP_CLAUSE_DEFAULTMAP:
case OMP_CLAUSE_AUTO:
case OMP_CLAUSE_SEQ:
+ case OMP_CLAUSE_NOHOST:
case OMP_CLAUSE_TILE:
WALK_SUBTREE_TAIL (OMP_CLAUSE_CHAIN (*tp));
@@ -1526,6 +1526,9 @@ extern void protected_set_expr_location (tree, location_t);
#define OMP_CLAUSE_VECTOR_LENGTH_EXPR(NODE) \
OMP_CLAUSE_OPERAND ( \
OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_VECTOR_LENGTH), 0)
+#define OMP_CLAUSE_BIND_NAME(NODE) \
+ OMP_CLAUSE_OPERAND ( \
+ OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_BIND), 0)
#define OMP_CLAUSE_DEPEND_KIND(NODE) \
(OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_DEPEND)->omp_clause.subcode.depend_kind)