OpenE2K
/
gcc
2
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

c-pragma.c (oacc_pragmas): Add entry for declare directive. 

2015-11-12  James Norris  <jnorris@codesourcery.com>
	    Joseph Myers  <joseph@codesourcery.com>

	gcc/c-family/
	* 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.

	gcc/c/
	* 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.

	gcc/cp/
	* 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.

	gcc/
	* 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.

	gcc/testsuite
	* c-c++-common/goacc/declare-1.c: New test.
	* c-c++-common/goacc/declare-2.c: Likewise.

	include/
	* gomp-constants.h (enum gomp_map_kind): Add GOMP_MAP_DEVICE_RESIDENT
	and GOMP_MAP_LINK.

	libgomp/
	* 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.

Co-Authored-By: Joseph Myers <joseph@codesourcery.com>

From-SVN: r230275
This commit is contained in:
James Norris 2015-11-12 22:20:41 +00:00 committed by James Norris
parent b6c34ca9ee
commit 6e232ba424
28 changed files with 1145 additions and 3 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

23
gcc/ChangeLog
View File

@ -1,3 +1,26 @@
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.

8
gcc/c-family/ChangeLog
View File

@ -1,3 +1,11 @@
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

1
gcc/c-family/c-pragma.c
View File

@ -1248,6 +1248,7 @@ 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 },
{ "kernels", PRAGMA_OACC_KERNELS },

5
gcc/c-family/c-pragma.h
View File

@ -30,6 +30,7 @@ enum pragma_kind {
PRAGMA_OACC_ATOMIC,
PRAGMA_OACC_CACHE,
PRAGMA_OACC_DATA,
PRAGMA_OACC_DECLARE,
PRAGMA_OACC_ENTER_DATA,
PRAGMA_OACC_EXIT_DATA,
PRAGMA_OACC_KERNELS,
@ -152,6 +153,7 @@ enum pragma_omp_clause {
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,
@ -176,7 +178,8 @@ enum pragma_omp_clause {
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;

12
gcc/c/ChangeLog
View File

@ -1,3 +1,15 @@
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

176
gcc/c/c-parser.c
View File

@ -1228,6 +1228,7 @@ static vec<tree, va_gc> *c_parser_expr_list (c_parser *, bool, bool,
vec<tree, va_gc> **, location_t *,
tree *, vec<location_t> *,
unsigned int * = NULL);
static 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 *);
@ -9729,6 +9730,10 @@ 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:
c_parser_oacc_enter_exit_data (parser, true);
return false;
@ -10018,6 +10023,8 @@ c_parser_omp_clause_name (c_parser *parser)
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;
@ -10454,10 +10461,16 @@ 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;
@ -12739,6 +12752,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
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";
@ -12761,6 +12778,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
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";
@ -13217,6 +13238,161 @@ c_parser_oacc_data (location_t loc, c_parser *parser)
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

14
gcc/cp/ChangeLog
View File

@ -1,3 +1,17 @@
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

172
gcc/cp/parser.c
View File

@ -29128,6 +29128,8 @@ cp_parser_omp_clause_name (cp_parser *parser)
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;
@ -29541,10 +29543,16 @@ 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;
@ -31545,6 +31553,10 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
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);
@ -31569,6 +31581,10 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
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";
@ -34525,6 +34541,158 @@ cp_parser_oacc_data (cp_parser *parser, cp_token *pragma_tok)
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
@ -36354,6 +36522,10 @@ cp_parser_pragma (cp_parser *parser, enum pragma_context context)
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;

8
gcc/cp/pt.c
View File

@ -15408,6 +15408,14 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
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:

3
gcc/gimple-pretty-print.c
View File

@ -1353,6 +1353,9 @@ dump_gimple_omp_target (pretty_printer *buffer, gomp_target *gs,
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 ();
}

2
gcc/gimple.h
View File

@ -170,6 +170,7 @@ enum gf_mask {
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
@ -6004,6 +6005,7 @@ is_gimple_omp_oacc (const gimple *stmt)
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;

184
gcc/gimplify.c
View File

@ -176,6 +176,7 @@ static struct gimplify_omp_ctx *gimplify_omp_ctxp;
/* 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. */
@ -1078,6 +1079,7 @@ gimplify_bind_expr (tree *expr_p, gimple_seq *pre_p)
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);
@ -1179,9 +1181,39 @@ gimplify_bind_expr (tree *expr_p, gimple_seq *pre_p)
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;
@ -5809,6 +5841,26 @@ omp_notice_threadprivate_variable (struct gimplify_omp_ctx *ctx, tree decl,
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.
@ -5908,6 +5960,15 @@ static unsigned
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)
{
@ -5928,7 +5989,7 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
|| 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
@ -7822,6 +7883,121 @@ gimplify_oacc_cache (tree *expr_p, gimple_seq *pre_p)
*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
@ -10182,11 +10358,15 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
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:

2
gcc/omp-builtins.def
View File

@ -353,3 +353,5 @@ DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_ENTER_EXIT_DATA,
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)

10
gcc/omp-low.c
View File

@ -12454,6 +12454,7 @@ expand_omp_target (struct omp_region *region)
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:
@ -12697,6 +12698,9 @@ expand_omp_target (struct omp_region *region)
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 ();
}
@ -12819,6 +12823,7 @@ expand_omp_target (struct omp_region *region)
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:
@ -13133,6 +13138,7 @@ build_omp_regions_1 (basic_block bb, struct omp_region *parent,
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;
@ -14916,6 +14922,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
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:
@ -14987,6 +14994,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
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:
@ -16713,6 +16722,7 @@ make_gimple_omp_edges (basic_block bb, struct omp_region **region,
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:

6
gcc/testsuite/ChangeLog
View File

@ -1,3 +1,9 @@
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.

83
gcc/testsuite/c-c++-common/goacc/declare-1.c Normal file
View File

@ -0,0 +1,83 @@
/* 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:;
}

79
gcc/testsuite/c-c++-common/goacc/declare-2.c Normal file
View File

@ -0,0 +1,79 @@
/* 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" } */
}

6
gcc/tree-pretty-print.c
View File

@ -654,6 +654,12 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, int flags)
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 ();
}

6
include/ChangeLog
View File

@ -1,3 +1,9 @@
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

5
include/gomp-constants.h
View File

@ -72,6 +72,11 @@ enum gomp_map_kind
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. */

11
libgomp/ChangeLog
View File

@ -1,3 +1,14 @@
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.

1
libgomp/libgomp.map
View File

@ -392,6 +392,7 @@ GOACC_2.0 {
GOACC_2.0.1 {
global:
GOACC_declare;
GOACC_parallel_keyed;
} GOACC_2.0;

58
libgomp/oacc-parallel.c
View File

@ -501,3 +501,61 @@ GOACC_get_thread_num (void)
{
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;
}
}
}

31
libgomp/testsuite/libgomp.oacc-c++/declare-1.C Normal file
View File

@ -0,0 +1,31 @@
/* { 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;
}

122
libgomp/testsuite/libgomp.oacc-c-c++-common/declare-1.c Normal file
View File

@ -0,0 +1,122 @@
/* { 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;
}

64
libgomp/testsuite/libgomp.oacc-c-c++-common/declare-2.c Normal file
View File

@ -0,0 +1,64 @@
/* { 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;
}

41
libgomp/testsuite/libgomp.oacc-c-c++-common/declare-4.c Normal file
View File

@ -0,0 +1,41 @@
/* { 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;
}

15
libgomp/testsuite/libgomp.oacc-c-c++-common/declare-5.c Normal file
View File

@ -0,0 +1,15 @@
/* { 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 "" } */