+2015-11-12 James Norris <jnorris@codesourcery.com>
+ Joseph Myers <joseph@codesourcery.com>
+
+ * gimple-pretty-print.c (dump_gimple_omp_target): Handle
+ GF_OMP_TARGET_KIND_OACC_DECLARE.
+ * gimple.h (enum gf_mask): Add GF_OMP_TARGET_KIND_OACC_DECLARE.
+ (is_gomple_omp_oacc): Handle GF_OMP_TARGET_KIND_OACC_DECLARE.
+ * gimplify.c (oacc_declare_returns): New.
+ (gimplify_bind_expr): Prepend 'exit' stmt to cleanup.
+ (device_resident_p): New function.
+ (oacc_default_clause): Handle device_resident clause.
+ (gimplify_oacc_declare_1, gimplify_oacc_declare): New functions.
+ (gimplify_expr): Handle OACC_DECLARE.
+ * omp-builtins.def (BUILT_IN_GOACC_DECLARE): New builtin.
+ * omp-low.c (expand_omp_target): Handle
+ GF_OMP_TARGET_KIND_OACC_DECLARE and BUILTIN_GOACC_DECLARE.
+ (build_omp_regions_1): Handlde GF_OMP_TARGET_KIND_OACC_DECLARE.
+ (lower_omp_target): Handle GF_OMP_TARGET_KIND_OACC_DECLARE,
+ GOMP_MAP_DEVICE_RESIDENT and GOMP_MAP_LINK.
+ (make_gimple_omp_edges): Handle GF_OMP_TARGET_KIND_OACC_DECLARE.
+ * tree-pretty-print.c (dump_omp_clause): Handle GOMP_MAP_LINK and
+ GOMP_MAP_DEVICE_RESIDENT.
+
2015-11-12 Christophe Lyon <christophe.lyon@linaro.org>
[ARM] Remove neon-testgen.ml and generated tests.
+2015-11-12 James Norris <jnorris@codesourcery.com>
+ Joseph Myers <joseph@codesourcery.com>
+
+ * c-pragma.c (oacc_pragmas): Add entry for declare directive.
+ * c-pragma.h (enum pragma_kind): Add PRAGMA_OACC_DECLARE.
+ (enum pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT and
+ PRAGMA_OACC_CLAUSE_LINK.
+
2015-11-11 Marek Polacek <polacek@redhat.com>
PR c/68107
{ "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 },
{ "kernels", PRAGMA_OACC_KERNELS },
PRAGMA_OACC_ATOMIC,
PRAGMA_OACC_CACHE,
PRAGMA_OACC_DATA,
+ PRAGMA_OACC_DECLARE,
PRAGMA_OACC_ENTER_DATA,
PRAGMA_OACC_EXIT_DATA,
PRAGMA_OACC_KERNELS,
PRAGMA_OACC_CLAUSE_CREATE,
PRAGMA_OACC_CLAUSE_DELETE,
PRAGMA_OACC_CLAUSE_DEVICEPTR,
+ PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT,
PRAGMA_OACC_CLAUSE_GANG,
PRAGMA_OACC_CLAUSE_HOST,
PRAGMA_OACC_CLAUSE_INDEPENDENT,
PRAGMA_OACC_CLAUSE_FIRSTPRIVATE = PRAGMA_OMP_CLAUSE_FIRSTPRIVATE,
PRAGMA_OACC_CLAUSE_IF = PRAGMA_OMP_CLAUSE_IF,
PRAGMA_OACC_CLAUSE_PRIVATE = PRAGMA_OMP_CLAUSE_PRIVATE,
- PRAGMA_OACC_CLAUSE_REDUCTION = PRAGMA_OMP_CLAUSE_REDUCTION
+ PRAGMA_OACC_CLAUSE_REDUCTION = PRAGMA_OMP_CLAUSE_REDUCTION,
+ PRAGMA_OACC_CLAUSE_LINK = PRAGMA_OMP_CLAUSE_LINK
};
extern struct cpp_reader* parse_in;
+2015-11-12 James Norris <jnorris@codesourcery.com>
+ Joseph Myers <joseph@codesourcery.com>
+
+ * c-parser.c (c_parser_pragma): Handle PRAGMA_OACC_DECLARE.
+ (c_parser_omp_clause_name): Handle 'device_resident' clause.
+ (c_parser_oacc_data_clause): Handle PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT
+ and PRAGMA_OMP_CLAUSE_LINK.
+ (c_parser_oacc_all_clauses): Handle PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT
+ and PRAGMA_OACC_CLAUSE_LINK.
+ (OACC_DECLARE_CLAUSE_MASK): New definition.
+ (c_parser_oacc_declare): New function.
+
2015-11-12 Marek Polacek <polacek@redhat.com>
PR c/67784
vec<tree, va_gc> **, location_t *,
tree *, vec<location_t> *,
unsigned int * = NULL);
+static void c_parser_oacc_declare (c_parser *);
static void c_parser_oacc_enter_exit_data (c_parser *, bool);
static void c_parser_oacc_update (c_parser *);
static void c_parser_omp_construct (c_parser *);
switch (id)
{
+ case PRAGMA_OACC_DECLARE:
+ c_parser_oacc_declare (parser);
+ return false;
+
case PRAGMA_OACC_ENTER_DATA:
c_parser_oacc_enter_exit_data (parser, true);
return false;
result = PRAGMA_OMP_CLAUSE_DEVICE;
else if (!strcmp ("deviceptr", p))
result = PRAGMA_OACC_CLAUSE_DEVICEPTR;
+ else if (!strcmp ("device_resident", p))
+ result = PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT;
else if (!strcmp ("dist_schedule", p))
result = PRAGMA_OMP_CLAUSE_DIST_SCHEDULE;
break;
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;
clauses = c_parser_oacc_data_clause_deviceptr (parser, clauses);
c_name = "deviceptr";
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_FIRSTPRIVATE:
clauses = c_parser_omp_clause_firstprivate (parser, clauses);
c_name = "firstprivate";
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_NUM_GANGS:
clauses = c_parser_omp_clause_num_gangs (parser, clauses);
c_name = "num_gangs";
return stmt;
}
+/* OpenACC 2.0:
+ # pragma acc declare oacc-data-clause[optseq] new-line
+*/
+
+#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_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) )
+
+static void
+c_parser_oacc_declare (c_parser *parser)
+{
+ location_t pragma_loc = c_parser_peek_token (parser)->location;
+ tree clauses, stmt, t, decl;
+
+ bool error = false;
+
+ 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 (t = clauses; t; t = OMP_CLAUSE_CHAIN (t))
+ {
+ location_t loc = OMP_CLAUSE_LOCATION (t);
+ decl = OMP_CLAUSE_DECL (t);
+ if (!DECL_P (decl))
+ {
+ error_at (loc, "array section in %<#pragma acc declare%>");
+ error = true;
+ continue;
+ }
+
+ 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 ()
+ && (TREE_STATIC (decl)
+ || !DECL_EXTERNAL (decl)))
+ {
+ error_at (loc,
+ "%qD must be a global variable in"
+ "%<#pragma acc declare link%>",
+ decl);
+ error = true;
+ continue;
+ }
+ break;
+
+ default:
+ if (global_bindings_p ())
+ {
+ error_at (loc, "invalid OpenACC clause at file scope");
+ error = true;
+ continue;
+ }
+ if (DECL_EXTERNAL (decl))
+ {
+ error_at (loc,
+ "invalid use of %<extern%> variable %qD "
+ "in %<#pragma acc declare%>", decl);
+ error = true;
+ continue;
+ }
+ else if (TREE_PUBLIC (decl))
+ {
+ error_at (loc,
+ "invalid use of %<global%> variable %qD "
+ "in %<#pragma acc declare%>", decl);
+ error = true;
+ continue;
+ }
+ break;
+ }
+
+ if (lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl))
+ || lookup_attribute ("omp declare target link",
+ DECL_ATTRIBUTES (decl)))
+ {
+ error_at (loc, "variable %qD used more than once with "
+ "%<#pragma acc declare%>", decl);
+ error = true;
+ continue;
+ }
+
+ if (!error)
+ {
+ tree id;
+
+ if (OMP_CLAUSE_MAP_KIND (t) == GOMP_MAP_LINK)
+ id = get_identifier ("omp declare target link");
+ else
+ id = get_identifier ("omp declare target");
+
+ DECL_ATTRIBUTES (decl)
+ = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (decl));
+
+ if (global_bindings_p ())
+ {
+ symtab_node *node = symtab_node::get (decl);
+ if (node != NULL)
+ {
+ node->offloadable = 1;
+#ifdef ENABLE_OFFLOADING
+ g->have_offload = true;
+ if (is_a <varpool_node *> (node))
+ {
+ vec_safe_push (offload_vars, decl);
+ node->force_output = 1;
+ }
+#endif
+ }
+ }
+ }
+ }
+
+ if (error || global_bindings_p ())
+ return;
+
+ stmt = make_node (OACC_DECLARE);
+ TREE_TYPE (stmt) = void_type_node;
+ OACC_DECLARE_CLAUSES (stmt) = clauses;
+ SET_EXPR_LOCATION (stmt, pragma_loc);
+
+ add_stmt (stmt);
+
+ return;
+}
+
/* OpenACC 2.0:
# pragma acc enter data oacc-enter-data-clause[optseq] new-line
+2015-11-12 James Norris <jnorris@codesourcery.com>
+ Joseph Myers <joseph@codesourcery.com>
+
+ * parser.c (cp_parser_omp_clause_name): Handle 'device_resident'
+ clause.
+ (cp_parser_oacc_data_clause): Handle PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT
+ and PRAGMA_OMP_CLAUSE_LINK.
+ (cp_paser_oacc_all_clauses): Handle PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT
+ and PRAGMA_OMP_CLAUSE_LINK.
+ (OACC_DECLARE_CLAUSE_MASK): New definition.
+ (cp_parser_oacc_declare): New function.
+ (cp_parser_pragma): Handle PRAGMA_OACC_DECLARE.
+ * pt.c (tsubst_expr): Handle OACC_DECLARE.
+
2015-11-12 Jason Merrill <jason@redhat.com>
* pt.c (check_explicit_specialization): Check the namespace after
result = PRAGMA_OMP_CLAUSE_DEVICE;
else if (!strcmp ("deviceptr", p))
result = PRAGMA_OACC_CLAUSE_DEVICEPTR;
+ else if (!strcmp ("device_resident", p))
+ result = PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT;
else if (!strcmp ("dist_schedule", p))
result = PRAGMA_OMP_CLAUSE_DIST_SCHEDULE;
break;
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;
clauses = cp_parser_oacc_data_clause_deviceptr (parser, clauses);
c_name = "deviceptr";
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_FIRSTPRIVATE:
clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_FIRSTPRIVATE,
clauses);
clauses, here);
c_name = "independent";
break;
+ case PRAGMA_OACC_CLAUSE_LINK:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "link";
+ break;
case PRAGMA_OACC_CLAUSE_NUM_GANGS:
code = OMP_CLAUSE_NUM_GANGS;
c_name = "num_gangs";
return stmt;
}
+/* OpenACC 2.0:
+ # pragma acc declare oacc-data-clause[optseq] new-line
+*/
+
+#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_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))
+
+static tree
+cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
+{
+ tree clauses, stmt, t;
+ bool error = false;
+
+ clauses = cp_parser_oacc_all_clauses (parser, OACC_DECLARE_CLAUSE_MASK,
+ "#pragma acc declare", pragma_tok, true);
+
+
+ if (find_omp_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE)
+ {
+ error_at (pragma_tok->location,
+ "no valid clauses specified in %<#pragma acc declare%>");
+ return NULL_TREE;
+ }
+
+ 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, "array section in %<#pragma acc declare%>");
+ error = true;
+ 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 ()
+ && (TREE_STATIC (decl)
+ || !DECL_EXTERNAL (decl)))
+ {
+ error_at (loc,
+ "%qD must be a global variable in"
+ "%<#pragma acc declare link%>",
+ decl);
+ error = true;
+ continue;
+ }
+ break;
+
+ default:
+ if (global_bindings_p ())
+ {
+ error_at (loc, "invalid OpenACC clause at file scope");
+ error = true;
+ continue;
+ }
+ if (DECL_EXTERNAL (decl))
+ {
+ error_at (loc,
+ "invalid use of %<extern%> variable %qD "
+ "in %<#pragma acc declare%>", decl);
+ error = true;
+ continue;
+ }
+ else if (TREE_PUBLIC (decl))
+ {
+ error_at (loc,
+ "invalid use of %<global%> variable %qD "
+ "in %<#pragma acc declare%>", decl);
+ error = true;
+ continue;
+ }
+ break;
+ }
+
+ if (lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl))
+ || lookup_attribute ("omp declare target link",
+ DECL_ATTRIBUTES (decl)))
+ {
+ error_at (loc, "variable %qD used more than once with "
+ "%<#pragma acc declare%>", decl);
+ error = true;
+ continue;
+ }
+
+ if (!error)
+ {
+ tree id;
+
+ if (OMP_CLAUSE_MAP_KIND (t) == GOMP_MAP_LINK)
+ id = get_identifier ("omp declare target link");
+ else
+ id = get_identifier ("omp declare target");
+
+ DECL_ATTRIBUTES (decl)
+ = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (decl));
+ if (global_bindings_p ())
+ {
+ symtab_node *node = symtab_node::get (decl);
+ if (node != NULL)
+ {
+ node->offloadable = 1;
+#ifdef ENABLE_OFFLOADING
+ g->have_offload = true;
+ if (is_a <varpool_node *> (node))
+ {
+ vec_safe_push (offload_vars, decl);
+ node->force_output = 1;
+ }
+#endif
+ }
+ }
+ }
+ }
+
+ if (error || global_bindings_p ())
+ return NULL_TREE;
+
+ stmt = make_node (OACC_DECLARE);
+ TREE_TYPE (stmt) = void_type_node;
+ OACC_DECLARE_CLAUSES (stmt) = clauses;
+ SET_EXPR_LOCATION (stmt, pragma_tok->location);
+
+ add_stmt (stmt);
+
+ return NULL_TREE;
+}
+
/* OpenACC 2.0:
# pragma acc enter data oacc-enter-data-clause[optseq] new-line
cp_parser_omp_declare (parser, pragma_tok, context);
return false;
+ case PRAGMA_OACC_DECLARE:
+ cp_parser_oacc_declare (parser, pragma_tok);
+ return false;
+
case PRAGMA_OACC_ROUTINE:
cp_parser_oacc_routine (parser, pragma_tok, context);
return false;
add_stmt (t);
break;
+ case OACC_DECLARE:
+ t = copy_node (t);
+ tmp = tsubst_omp_clauses (OACC_DECLARE_CLAUSES (t), false, false,
+ args, complain, in_decl);
+ OACC_DECLARE_CLAUSES (t) = tmp;
+ add_stmt (t);
+ break;
+
case OMP_TARGET_UPDATE:
case OMP_TARGET_ENTER_DATA:
case OMP_TARGET_EXIT_DATA:
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
kind = " oacc_enter_exit_data";
break;
+ case GF_OMP_TARGET_KIND_OACC_DECLARE:
+ kind = " oacc_declare";
+ break;
default:
gcc_unreachable ();
}
GF_OMP_TARGET_KIND_OACC_DATA = 7,
GF_OMP_TARGET_KIND_OACC_UPDATE = 8,
GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA = 9,
+ GF_OMP_TARGET_KIND_OACC_DECLARE = 10,
/* True on an GIMPLE_OMP_RETURN statement if the return does not require
a thread synchronization via some sort of barrier. The exact barrier
case GF_OMP_TARGET_KIND_OACC_DATA:
case GF_OMP_TARGET_KIND_OACC_UPDATE:
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ case GF_OMP_TARGET_KIND_OACC_DECLARE:
return true;
default:
return false;
/* Forward declaration. */
static enum gimplify_status gimplify_compound_expr (tree *, gimple_seq *, bool);
+static hash_map<tree, tree> *oacc_declare_returns;
/* Shorter alias name for the above function for use in gimplify.c
only. */
gimple_seq body, cleanup;
gcall *stack_save;
location_t start_locus = 0, end_locus = 0;
+ tree ret_clauses = NULL;
tree temp = voidify_wrapper_expr (bind_expr, NULL);
clobber_stmt = gimple_build_assign (t, clobber);
gimple_set_location (clobber_stmt, end_locus);
gimplify_seq_add_stmt (&cleanup, clobber_stmt);
+
+ if (flag_openacc && oacc_declare_returns != NULL)
+ {
+ tree *c = oacc_declare_returns->get (t);
+ if (c != NULL)
+ {
+ if (ret_clauses)
+ OMP_CLAUSE_CHAIN (*c) = ret_clauses;
+
+ ret_clauses = *c;
+
+ oacc_declare_returns->remove (t);
+
+ if (oacc_declare_returns->elements () == 0)
+ {
+ delete oacc_declare_returns;
+ oacc_declare_returns = NULL;
+ }
+ }
+ }
}
}
+ if (ret_clauses)
+ {
+ gomp_target *stmt;
+ gimple_stmt_iterator si = gsi_start (cleanup);
+
+ stmt = gimple_build_omp_target (NULL, GF_OMP_TARGET_KIND_OACC_DECLARE,
+ ret_clauses);
+ gsi_insert_seq_before_without_update (&si, stmt, GSI_NEW_STMT);
+ }
+
if (cleanup)
{
gtry *gs;
return false;
}
+/* Return true if global var DECL is device resident. */
+
+static bool
+device_resident_p (tree decl)
+{
+ tree attr = lookup_attribute ("oacc declare target", DECL_ATTRIBUTES (decl));
+
+ if (!attr)
+ return false;
+
+ for (tree t = TREE_VALUE (attr); t; t = TREE_PURPOSE (t))
+ {
+ tree c = TREE_VALUE (t);
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DEVICE_RESIDENT)
+ return true;
+ }
+
+ return false;
+}
+
/* Determine outer default flags for DECL mentioned in an OMP region
but not declared in an enclosing clause.
oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
{
const char *rkind;
+ bool on_device = false;
+
+ if ((ctx->region_type & (ORT_ACC_PARALLEL | ORT_ACC_KERNELS)) != 0
+ && is_global_var (decl)
+ && device_resident_p (decl))
+ {
+ on_device = true;
+ flags |= GOVD_MAP_TO_ONLY;
+ }
switch (ctx->region_type)
{
|| POINTER_TYPE_P (type))
type = TREE_TYPE (type);
- if (AGGREGATE_TYPE_P (type))
+ if (on_device || AGGREGATE_TYPE_P (type))
/* Aggregates default to 'present_or_copy'. */
flags |= GOVD_MAP;
else
*expr_p = NULL_TREE;
}
+/* Helper function of gimplify_oacc_declare. The helper's purpose is to,
+ if required, translate 'kind' in CLAUSE into an 'entry' kind and 'exit'
+ kind. The entry kind will replace the one in CLAUSE, while the exit
+ kind will be used in a new omp_clause and returned to the caller. */
+
+static tree
+gimplify_oacc_declare_1 (tree clause)
+{
+ HOST_WIDE_INT kind, new_op;
+ bool ret = false;
+ tree c = NULL;
+
+ kind = OMP_CLAUSE_MAP_KIND (clause);
+
+ switch (kind)
+ {
+ case GOMP_MAP_ALLOC:
+ case GOMP_MAP_FORCE_ALLOC:
+ case GOMP_MAP_FORCE_TO:
+ new_op = GOMP_MAP_FORCE_DEALLOC;
+ ret = true;
+ break;
+
+ case GOMP_MAP_FORCE_FROM:
+ OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_FORCE_ALLOC);
+ new_op = GOMP_MAP_FORCE_FROM;
+ ret = true;
+ break;
+
+ case GOMP_MAP_FORCE_TOFROM:
+ OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_FORCE_TO);
+ new_op = GOMP_MAP_FORCE_FROM;
+ ret = true;
+ break;
+
+ case GOMP_MAP_FROM:
+ OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_FORCE_ALLOC);
+ new_op = GOMP_MAP_FROM;
+ ret = true;
+ break;
+
+ case GOMP_MAP_TOFROM:
+ OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_TO);
+ new_op = GOMP_MAP_FROM;
+ ret = true;
+ break;
+
+ case GOMP_MAP_DEVICE_RESIDENT:
+ case GOMP_MAP_FORCE_DEVICEPTR:
+ case GOMP_MAP_FORCE_PRESENT:
+ case GOMP_MAP_LINK:
+ case GOMP_MAP_POINTER:
+ case GOMP_MAP_TO:
+ break;
+
+ default:
+ gcc_unreachable ();
+ break;
+ }
+
+ if (ret)
+ {
+ c = build_omp_clause (OMP_CLAUSE_LOCATION (clause), OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (c, new_op);
+ OMP_CLAUSE_DECL (c) = OMP_CLAUSE_DECL (clause);
+ }
+
+ return c;
+}
+
+/* Gimplify OACC_DECLARE. */
+
+static void
+gimplify_oacc_declare (tree *expr_p, gimple_seq *pre_p)
+{
+ tree expr = *expr_p;
+ gomp_target *stmt;
+ tree clauses, t;
+
+ clauses = OACC_DECLARE_CLAUSES (expr);
+
+ gimplify_scan_omp_clauses (&clauses, pre_p, ORT_TARGET_DATA, OACC_DECLARE);
+
+ for (t = clauses; t; t = OMP_CLAUSE_CHAIN (t))
+ {
+ tree decl = OMP_CLAUSE_DECL (t);
+
+ if (TREE_CODE (decl) == MEM_REF)
+ continue;
+
+ if (TREE_CODE (decl) == VAR_DECL
+ && !is_global_var (decl)
+ && DECL_CONTEXT (decl) == current_function_decl)
+ {
+ tree c = gimplify_oacc_declare_1 (t);
+ if (c)
+ {
+ if (oacc_declare_returns == NULL)
+ oacc_declare_returns = new hash_map<tree, tree>;
+
+ oacc_declare_returns->put (decl, c);
+ }
+ }
+
+ omp_add_variable (gimplify_omp_ctxp, decl, GOVD_SEEN);
+ }
+
+ stmt = gimple_build_omp_target (NULL, GF_OMP_TARGET_KIND_OACC_DECLARE,
+ clauses);
+
+ gimplify_seq_add_stmt (pre_p, stmt);
+
+ *expr_p = NULL_TREE;
+}
+
/* Gimplify the contents of an OMP_PARALLEL statement. This involves
gimplification of the body, as well as scanning the body for used
variables. We need to do this scan now, because variable-sized
break;
case OACC_HOST_DATA:
- case OACC_DECLARE:
sorry ("directive not yet implemented");
ret = GS_ALL_DONE;
break;
+ case OACC_DECLARE:
+ gimplify_oacc_declare (expr_p, pre_p);
+ ret = GS_ALL_DONE;
+ break;
+
case OACC_DATA:
case OACC_KERNELS:
case OACC_PARALLEL:
BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_UINT_PTR, ATTR_NOTHROW_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TEAMS, "GOMP_teams",
BT_FN_VOID_UINT_UINT, ATTR_NOTHROW_LIST)
+DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DECLARE, "GOACC_declare",
+ BT_FN_VOID_INT_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
case GF_OMP_TARGET_KIND_OACC_KERNELS:
case GF_OMP_TARGET_KIND_OACC_UPDATE:
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ case GF_OMP_TARGET_KIND_OACC_DECLARE:
data_region = false;
break;
case GF_OMP_TARGET_KIND_DATA:
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
start_ix = BUILT_IN_GOACC_ENTER_EXIT_DATA;
break;
+ case GF_OMP_TARGET_KIND_OACC_DECLARE:
+ start_ix = BUILT_IN_GOACC_DECLARE;
+ break;
default:
gcc_unreachable ();
}
switch (start_ix)
{
case BUILT_IN_GOACC_DATA_START:
+ case BUILT_IN_GOACC_DECLARE:
case BUILT_IN_GOMP_TARGET_DATA:
break;
case BUILT_IN_GOMP_TARGET:
case GF_OMP_TARGET_KIND_EXIT_DATA:
case GF_OMP_TARGET_KIND_OACC_UPDATE:
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ case GF_OMP_TARGET_KIND_OACC_DECLARE:
/* ..., other than for those stand-alone directives... */
region = NULL;
break;
case GF_OMP_TARGET_KIND_OACC_KERNELS:
case GF_OMP_TARGET_KIND_OACC_UPDATE:
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ case GF_OMP_TARGET_KIND_OACC_DECLARE:
data_region = false;
break;
case GF_OMP_TARGET_KIND_DATA:
case GOMP_MAP_FORCE_TOFROM:
case GOMP_MAP_FORCE_PRESENT:
case GOMP_MAP_FORCE_DEVICEPTR:
+ case GOMP_MAP_DEVICE_RESIDENT:
+ case GOMP_MAP_LINK:
gcc_assert (is_gimple_omp_oacc (stmt));
break;
default:
case GF_OMP_TARGET_KIND_EXIT_DATA:
case GF_OMP_TARGET_KIND_OACC_UPDATE:
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ case GF_OMP_TARGET_KIND_OACC_DECLARE:
cur_region = cur_region->outer;
break;
default:
+2015-11-12 James Norris <jnorris@codesourcery.com>
+ Joseph Myers <joseph@codesourcery.com>
+
+ * c-c++-common/goacc/declare-1.c: New test.
+ * c-c++-common/goacc/declare-2.c: Likewise.
+
2015-11-12 Christophe Lyon <christophe.lyon@linaro.org>
[ARM] Remove neon-testgen.ml and generated tests.
--- /dev/null
+/* Test valid uses of declare directive. */
+/* { dg-do compile } */
+
+int v0;
+#pragma acc declare create(v0)
+
+int v1;
+#pragma acc declare copyin(v1)
+
+int *v2;
+#pragma acc declare deviceptr(v2)
+
+int v3;
+#pragma acc declare device_resident(v3)
+
+int v4;
+#pragma acc declare link(v4)
+
+int v5, v6, v7, v8;
+#pragma acc declare create(v5, v6) copyin(v7, v8)
+
+void
+f (void)
+{
+ int va0;
+#pragma acc declare create(va0)
+
+ int va1;
+#pragma acc declare copyin(va1)
+
+ int *va2;
+#pragma acc declare deviceptr(va2)
+
+ int va3;
+#pragma acc declare device_resident(va3)
+
+ extern int ve0;
+#pragma acc declare create(ve0)
+
+ extern int ve1;
+#pragma acc declare copyin(ve1)
+
+ extern int *ve2;
+#pragma acc declare deviceptr(ve2)
+
+ extern int ve3;
+#pragma acc declare device_resident(ve3)
+
+ extern int ve4;
+#pragma acc declare link(ve4)
+
+ int va5;
+#pragma acc declare copy(va5)
+
+ int va6;
+#pragma acc declare copyout(va6)
+
+ int va7;
+#pragma acc declare present(va7)
+
+ int va8;
+#pragma acc declare present_or_copy(va8)
+
+ int va9;
+#pragma acc declare present_or_copyin(va9)
+
+ int va10;
+#pragma acc declare present_or_copyout(va10)
+
+ int va11;
+#pragma acc declare present_or_create(va11)
+
+ a:
+ {
+ int va0;
+#pragma acc declare create(va0)
+ if (v1)
+ goto a;
+ else
+ goto b;
+ }
+ b:;
+}
--- /dev/null
+/* Test invalid uses of declare directive. */
+/* { dg-do compile } */
+
+#pragma acc declare /* { dg-error "no valid clauses" } */
+
+#pragma acc declare create(undeclared) /* { dg-error "undeclared" } */
+/* { dg-error "no valid clauses" "second error" { target *-*-* } 6 } */
+
+int v0[10];
+#pragma acc declare create(v0[1:3]) /* { dg-error "array section" } */
+
+int v1;
+#pragma acc declare create(v1, v1) /* { dg-error "more than once" } */
+
+int v2;
+#pragma acc declare create(v2)
+#pragma acc declare copyin(v2) /* { dg-error "more than once" } */
+
+int v3;
+#pragma acc declare copy(v3) /* { dg-error "at file scope" } */
+
+int v4;
+#pragma acc declare copyout(v4) /* { dg-error "at file scope" } */
+
+int v5;
+#pragma acc declare present(v5) /* { dg-error "at file scope" } */
+
+int v6;
+#pragma acc declare present_or_copy(v6) /* { dg-error "at file scope" } */
+
+int v7;
+#pragma acc declare present_or_copyin(v7) /* { dg-error "at file scope" } */
+
+int v8;
+#pragma acc declare present_or_copyout(v8) /* { dg-error "at file scope" } */
+
+int v9;
+#pragma acc declare present_or_create(v9) /* { dg-error "at file scope" } */
+
+int va10;
+#pragma acc declare create (va10)
+#pragma acc declare link (va10) /* { dg-error "more than once" } */
+
+int va11;
+#pragma acc declare link (va11)
+#pragma acc declare link (va11) /* { dg-error "more than once" } */
+
+int va12;
+#pragma acc declare create (va12) link (va12) /* { dg-error "more than once" } */
+
+void
+f (void)
+{
+ int va0;
+#pragma acc declare link(va0) /* { dg-error "global variable" } */
+
+ extern int ve0;
+#pragma acc declare copy(ve0) /* { dg-error "invalid use of" } */
+
+ extern int ve1;
+#pragma acc declare copyout(ve1) /* { dg-error "invalid use of" } */
+
+ extern int ve2;
+#pragma acc declare present(ve2) /* { dg-error "invalid use of" } */
+
+ extern int ve3;
+#pragma acc declare present_or_copy(ve3) /* { dg-error "invalid use of" } */
+
+ extern int ve4;
+#pragma acc declare present_or_copyin(ve4) /* { dg-error "invalid use of" } */
+
+ extern int ve5;
+#pragma acc declare present_or_copyout(ve5) /* { dg-error "invalid use of" } */
+
+ extern int ve6;
+#pragma acc declare present_or_create(ve6) /* { dg-error "invalid use of" } */
+
+#pragma acc declare present (v9) /* { dg-error "invalid use of" } */
+}
case GOMP_MAP_ALWAYS_POINTER:
pp_string (pp, "always_pointer");
break;
+ case GOMP_MAP_DEVICE_RESIDENT:
+ pp_string (pp, "device_resident");
+ break;
+ case GOMP_MAP_LINK:
+ pp_string (pp, "link");
+ break;
default:
gcc_unreachable ();
}
+2015-11-12 James Norris <jnorris@codesourcery.com>
+ Joseph Myers <joseph@codesourcery.com>
+
+ * gomp-constants.h (enum gomp_map_kind): Add GOMP_MAP_DEVICE_RESIDENT
+ and GOMP_MAP_LINK.
+
2015-11-09 Alan Modra <amodra@gmail.com>
PR gdb/17133
POINTER_SIZE_UNITS. */
GOMP_MAP_FORCE_DEVICEPTR = (GOMP_MAP_FLAG_SPECIAL_1 | 0),
/* Do not map, copy bits for firstprivate instead. */
+ /* OpenACC device_resident. */
+ GOMP_MAP_DEVICE_RESIDENT = (GOMP_MAP_FLAG_SPECIAL_1 | 1),
+ /* OpenACC link. */
+ GOMP_MAP_LINK = (GOMP_MAP_FLAG_SPECIAL_1 | 2),
+ /* Allocate. */
GOMP_MAP_FIRSTPRIVATE = (GOMP_MAP_FLAG_SPECIAL | 0),
/* Similarly, but store the value in the pointer rather than
pointed by the pointer. */
+2015-11-12 James Norris <jnorris@codesourcery.com>
+ Joseph Myers <joseph@codesourcery.com>
+
+ * libgomp.map (GOACC_2.0.1): Export GOACC_declare.
+ * oacc-parallel.c (GOACC_declare): New function.
+ * testsuite/libgomp.oacc-c-c++-common/declare-1.c: New test.
+ * testsuite/libgomp.oacc-c-c++-common/declare-2.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/declare-4.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/declare-5.c: Likewise.
+ * testsuite/libgomp.oacc-c++/declare-1.C: Likewise.
+
2015-11-12 Nathan Sidwell <nathan@codesourcery.com>
* testsuite/libgomp.oacc-c-c++-common/default-1.c: New.
GOACC_2.0.1 {
global:
+ GOACC_declare;
GOACC_parallel_keyed;
} GOACC_2.0;
{
return 0;
}
+
+void
+GOACC_declare (int device, size_t mapnum,
+ void **hostaddrs, size_t *sizes, unsigned short *kinds)
+{
+ int i;
+
+ for (i = 0; i < mapnum; i++)
+ {
+ unsigned char kind = kinds[i] & 0xff;
+
+ if (kind == GOMP_MAP_POINTER || kind == GOMP_MAP_TO_PSET)
+ continue;
+
+ switch (kind)
+ {
+ case GOMP_MAP_FORCE_ALLOC:
+ case GOMP_MAP_FORCE_DEALLOC:
+ case GOMP_MAP_FORCE_FROM:
+ case GOMP_MAP_FORCE_TO:
+ case GOMP_MAP_POINTER:
+ GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i],
+ &kinds[i], 0, 0);
+ break;
+
+ case GOMP_MAP_FORCE_DEVICEPTR:
+ break;
+
+ case GOMP_MAP_ALLOC:
+ if (!acc_is_present (hostaddrs[i], sizes[i]))
+ GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i],
+ &kinds[i], 0, 0);
+ break;
+
+ case GOMP_MAP_TO:
+ GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i],
+ &kinds[i], 0, 0);
+
+ break;
+
+ case GOMP_MAP_FROM:
+ kinds[i] = GOMP_MAP_FORCE_FROM;
+ GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i],
+ &kinds[i], 0, 0);
+ break;
+
+ case GOMP_MAP_FORCE_PRESENT:
+ if (!acc_is_present (hostaddrs[i], sizes[i]))
+ gomp_fatal ("[%p,%ld] is not mapped", hostaddrs[i],
+ (unsigned long) sizes[i]);
+ break;
+
+ default:
+ assert (0);
+ break;
+ }
+ }
+}
--- /dev/null
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
+
+#include <stdlib.h>
+
+template<class T>
+T foo()
+{
+ T a, b;
+ #pragma acc declare create (a)
+
+ #pragma acc parallel copyout (b)
+ {
+ a = 5;
+ b = a;
+ }
+
+ return b;
+}
+
+int
+main (void)
+{
+ int rc;
+
+ rc = foo<int>();
+
+ if (rc != 5)
+ abort ();
+
+ return 0;
+}
--- /dev/null
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
+
+#include <openacc.h>
+#include <stdlib.h>
+#include <stdio.h>
+
+#define N 8
+
+void
+subr2 (int *a)
+{
+ int i;
+ int f[N];
+#pragma acc declare copyout (f)
+
+#pragma acc parallel copy (a[0:N])
+ {
+ for (i = 0; i < N; i++)
+ {
+ f[i] = a[i];
+ a[i] = f[i] + f[i] + f[i];
+ }
+ }
+}
+
+void
+subr1 (int *a)
+{
+ int f[N];
+#pragma acc declare copy (f)
+
+#pragma acc parallel copy (a[0:N])
+ {
+ int i;
+
+ for (i = 0; i < N; i++)
+ {
+ f[i] = a[i];
+ a[i] = f[i] + f[i];
+ }
+ }
+}
+
+int b[8];
+#pragma acc declare create (b)
+
+int d[8] = { 1, 2, 3, 4, 5, 6, 7, 8 };
+#pragma acc declare copyin (d)
+
+int
+main (int argc, char **argv)
+{
+ int a[N];
+ int e[N];
+#pragma acc declare create (e)
+ int i;
+
+ for (i = 0; i < N; i++)
+ a[i] = i + 1;
+
+ if (!acc_is_present (&b, sizeof (b)))
+ abort ();
+
+ if (!acc_is_present (&d, sizeof (d)))
+ abort ();
+
+ if (!acc_is_present (&e, sizeof (e)))
+ abort ();
+
+#pragma acc parallel copyin (a[0:N])
+ {
+ for (i = 0; i < N; i++)
+ {
+ b[i] = a[i];
+ a[i] = b[i];
+ }
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != i + 1)
+ abort ();
+ }
+
+#pragma acc parallel copy (a[0:N])
+ {
+ for (i = 0; i < N; i++)
+ {
+ e[i] = a[i] + d[i];
+ a[i] = e[i];
+ }
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != (i + 1) * 2)
+ abort ();
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 1234;
+ }
+
+ subr1 (&a[0]);
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 1234 * 2)
+ abort ();
+ }
+
+ subr2 (&a[0]);
+
+ for (i = 0; i < 1; i++)
+ {
+ if (a[i] != 1234 * 6)
+ abort ();
+ }
+
+ return 0;
+}
--- /dev/null
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
+
+#include <stdlib.h>
+
+#define N 16
+
+float c[N];
+#pragma acc declare device_resident (c)
+
+#pragma acc routine
+float
+subr2 (float a)
+{
+ int i;
+
+ for (i = 0; i < N; i++)
+ c[i] = 2.0;
+
+ for (i = 0; i < N; i++)
+ a += c[i];
+
+ return a;
+}
+
+float b[N];
+#pragma acc declare copyin (b)
+
+#pragma acc routine
+float
+subr1 (float a)
+{
+ int i;
+
+ for (i = 0; i < N; i++)
+ a += b[i];
+
+ return a;
+}
+
+int
+main (int argc, char **argv)
+{
+ float a;
+ int i;
+
+ for (i = 0; i < 16; i++)
+ b[i] = 1.0;
+
+ a = 0.0;
+
+ a = subr1 (a);
+
+ if (a != 16.0)
+ abort ();
+
+ a = 0.0;
+
+ a = subr2 (a);
+
+ if (a != 32.0)
+ abort ();
+
+ return 0;
+}
--- /dev/null
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
+
+#include <stdlib.h>
+#include <openacc.h>
+
+float b;
+#pragma acc declare link (b)
+
+#pragma acc routine
+int
+func (int a)
+{
+ b = a + 1;
+
+ return b;
+}
+
+int
+main (int argc, char **argv)
+{
+ float a;
+
+ a = 2.0;
+
+#pragma acc parallel copy (a)
+ {
+ b = a;
+ a = 1.0;
+ a = a + b;
+ }
+
+ if (a != 3.0)
+ abort ();
+
+ a = func (a);
+
+ if (a != 4.0)
+ abort ();
+
+ return 0;
+}
--- /dev/null
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
+
+#include <stdio.h>
+
+int
+main (int argc, char **argv)
+{
+ int a[8] __attribute__((unused));
+
+ fprintf (stderr, "CheCKpOInT\n");
+#pragma acc declare present (a)
+}
+
+/* { dg-output "CheCKpOInT" } */
+/* { dg-shouldfail "" } */
projects / thirdparty / gcc.git / commitdiff
