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

OpenACC host_data support.

Browse Source 
gcc/
	* gimple-pretty-print.c (dump_gimple_omp_target): Add host_data
	support.
	* gimple.h (gf_mask): Add GF_OMP_TARGET_KIND_OACC_HOST_DATA.
	(is_gimple_omp_oacc): Add support for above.
	* gimplify.c (omp_region_type): Add ORT_ACC_HOST_DATA.
	(omp_notice_variable): Diagnose undefined implicit uses of
	use_device variables in offloaded regions.
	(gimplify_scan_omp_clauses): Add host_data, use_device
	support. Diagnose undefined mapping of use_device variables in
	OpenACC clauses.
	(gimplify_omp_workshare): Add host_data support.
	(gimplify_expr): Likewise.
	* omp-builtins.def (BUILT_IN_GOACC_HOST_DATA): New.
	* omp-low.c (lookup_decl_in_outer_ctx)
	(maybe_lookup_decl_in_outer_ctx): Add optional argument to skip
	host_data regions.
	(scan_sharing_clauses): Support use_device.
	(check_omp_nesting_restrictions): Support host_data.
	(expand_omp_target): Support host_data.
	(lower_omp_target): Skip over outer host_data regions when looking
	up decls. Support use_device.
	(make_gimple_omp_edges): Support host_data.
	* tree-nested.c (convert_nonlocal_omp_clauses): Add use_device
	clause.

	gcc/c/
	* c-parser.c (c_parser_omp_clause_name): Add use_device support.
	(c_parser_oacc_clause_use_device): New function.
	(c_parser_oacc_all_clauses): Add use_device support.
	(OACC_HOST_DATA_CLAUSE_MASK): New macro.
	(c_parser_oacc_host_data): New function.
	(c_parser_omp_construct): Add host_data support.
	* c-tree.h (c_finish_oacc_host_data): Add prototype.
	* c-typeck.c (c_finish_oacc_host_data): New function.
	(c_finish_omp_clauses): Add use_device support.

	gcc/cp/
	* cp-tree.h (finish_oacc_host_data): Add prototype.
	* parser.c (cp_parser_omp_clause_name): Add use_device support.
	(cp_parser_oacc_all_clauses): Add use_device support.
	(OACC_HOST_DATA_CLAUSE_MASK): New macro.
	(cp_parser_oacc_host_data): New function.
	(cp_parser_omp_construct): Add host_data support.
	(cp_parser_pragma): Add host_data support.
	* semantics.c (finish_omp_clauses): Add use_device support.
	(finish_oacc_host_data): New function.

	gcc/c-family/
	* c-pragma.c (oacc_pragmas): Add PRAGMA_OACC_HOST_DATA.
	* c-pragma.h (pragma_kind): Add PRAGMA_OACC_HOST_DATA.
	(pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_USE_DEVICE.

	libgomp/
	* oacc-parallel.c (GOACC_host_data): New function.
	* libgomp.map (GOACC_host_data): Add to GOACC_2.0.1.
	* testsuite/libgomp.oacc-c-c++-common/host_data-1.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/host_data-2.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/host_data-3.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/host_data-4.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/host_data-5.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/host_data-6.c: New test.


Co-Authored-By: Cesar Philippidis <cesar@codesourcery.com>
Co-Authored-By: James Norris <James_Norris@mentor.com>

From-SVN: r231118
This commit is contained in:
Julian Brown 2015-12-01 15:52:23 +00:00 committed by Julian Brown
parent 4bc84763c0
commit 37d5ad46dd
27 changed files with 558 additions and 10 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

29
gcc/ChangeLog
View File

@ -1,3 +1,32 @@
2015-12-01 Julian Brown <julian@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
James Norris <James_Norris@mentor.com>
* gimple-pretty-print.c (dump_gimple_omp_target): Add host_data
support.
* gimple.h (gf_mask): Add GF_OMP_TARGET_KIND_OACC_HOST_DATA.
(is_gimple_omp_oacc): Add support for above.
* gimplify.c (omp_region_type): Add ORT_ACC_HOST_DATA.
(omp_notice_variable): Diagnose undefined implicit uses of
use_device variables in offloaded regions.
(gimplify_scan_omp_clauses): Add host_data, use_device
support. Diagnose undefined mapping of use_device variables in
OpenACC clauses.
(gimplify_omp_workshare): Add host_data support.
(gimplify_expr): Likewise.
* omp-builtins.def (BUILT_IN_GOACC_HOST_DATA): New.
* omp-low.c (lookup_decl_in_outer_ctx)
(maybe_lookup_decl_in_outer_ctx): Add optional argument to skip
host_data regions.
(scan_sharing_clauses): Support use_device.
(check_omp_nesting_restrictions): Support host_data.
(expand_omp_target): Support host_data.
(lower_omp_target): Skip over outer host_data regions when looking
up decls. Support use_device.
(make_gimple_omp_edges): Support host_data.
* tree-nested.c (convert_nonlocal_omp_clauses): Add use_device
clause.
2015-12-01 Marek Polacek <polacek@redhat.com>
PR middle-end/68582

8
gcc/c-family/ChangeLog
View File

@ -1,3 +1,11 @@
2015-12-01 Julian Brown <julian@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
James Norris <James_Norris@mentor.com>
* c-pragma.c (oacc_pragmas): Add PRAGMA_OACC_HOST_DATA.
* c-pragma.h (pragma_kind): Add PRAGMA_OACC_HOST_DATA.
(pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_USE_DEVICE.
2015-11-30 Eric Botcazou <ebotcazou@adacore.com>
* c-ada-spec.c (print_ada_macros): Remove redundant blank line.

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

@ -1251,6 +1251,7 @@ static const struct omp_pragma_def oacc_pragmas[] = {
{ "declare", PRAGMA_OACC_DECLARE },
{ "enter", PRAGMA_OACC_ENTER_DATA },
{ "exit", PRAGMA_OACC_EXIT_DATA },
{ "host_data", PRAGMA_OACC_HOST_DATA },
{ "kernels", PRAGMA_OACC_KERNELS },
{ "loop", PRAGMA_OACC_LOOP },
{ "parallel", PRAGMA_OACC_PARALLEL },

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

@ -33,6 +33,7 @@ enum pragma_kind {
PRAGMA_OACC_DECLARE,
PRAGMA_OACC_ENTER_DATA,
PRAGMA_OACC_EXIT_DATA,
PRAGMA_OACC_HOST_DATA,
PRAGMA_OACC_KERNELS,
PRAGMA_OACC_LOOP,
PRAGMA_OACC_PARALLEL,
@ -167,6 +168,7 @@ enum pragma_omp_clause {
PRAGMA_OACC_CLAUSE_SELF,
PRAGMA_OACC_CLAUSE_SEQ,
PRAGMA_OACC_CLAUSE_TILE,
PRAGMA_OACC_CLAUSE_USE_DEVICE,
PRAGMA_OACC_CLAUSE_VECTOR,
PRAGMA_OACC_CLAUSE_VECTOR_LENGTH,
PRAGMA_OACC_CLAUSE_WAIT,

14
gcc/c/ChangeLog
View File

@ -1,3 +1,17 @@
2015-12-01 Julian Brown <julian@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
James Norris <James_Norris@mentor.com>
* c-parser.c (c_parser_omp_clause_name): Add use_device support.
(c_parser_oacc_clause_use_device): New function.
(c_parser_oacc_all_clauses): Add use_device support.
(OACC_HOST_DATA_CLAUSE_MASK): New macro.
(c_parser_oacc_host_data): New function.
(c_parser_omp_construct): Add host_data support.
* c-tree.h (c_finish_oacc_host_data): Add prototype.
* c-typeck.c (c_finish_oacc_host_data): New function.
(c_finish_omp_clauses): Add use_device support.
2015-11-29 Jan Hubicka <hubicka@ucw.cz>
PR c/67106

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

@ -10279,6 +10279,8 @@ c_parser_omp_clause_name (c_parser *parser)
result = PRAGMA_OMP_CLAUSE_UNTIED;
else if (!strcmp ("use_device_ptr", p))
result = PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR;
else if (!strcmp ("use_device", p))
result = PRAGMA_OACC_CLAUSE_USE_DEVICE;
break;
case 'v':
if (!strcmp ("vector", p))
@ -11631,6 +11633,15 @@ c_parser_oacc_clause_tile (c_parser *parser, tree list)
return c;
}
/* OpenACC 2.0:
use_device ( variable-list ) */
static tree
c_parser_oacc_clause_use_device (c_parser *parser, tree list)
{
return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_USE_DEVICE, list);
}
/* OpenACC:
wait ( int-expr-list ) */
@ -12940,6 +12951,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "self";
break;
case PRAGMA_OACC_CLAUSE_USE_DEVICE:
clauses = c_parser_oacc_clause_use_device (parser, clauses);
c_name = "use_device";
break;
case PRAGMA_OACC_CLAUSE_SEQ:
clauses = c_parser_oacc_simple_clause (parser, OMP_CLAUSE_SEQ,
clauses);
@ -13589,6 +13604,29 @@ c_parser_oacc_enter_exit_data (c_parser *parser, bool enter)
}
/* OpenACC 2.0:
# pragma acc host_data oacc-data-clause[optseq] new-line
structured-block
*/
#define OACC_HOST_DATA_CLAUSE_MASK \
( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_USE_DEVICE) )
static tree
c_parser_oacc_host_data (location_t loc, c_parser *parser)
{
tree stmt, clauses, block;
clauses = c_parser_oacc_all_clauses (parser, OACC_HOST_DATA_CLAUSE_MASK,
"#pragma acc host_data");
block = c_begin_omp_parallel ();
add_stmt (c_parser_omp_structured_block (parser));
stmt = c_finish_oacc_host_data (loc, clauses, block);
return stmt;
}
/* OpenACC 2.0:
# pragma acc loop oacc-loop-clause[optseq] new-line
@ -16897,6 +16935,9 @@ c_parser_omp_construct (c_parser *parser)
case PRAGMA_OACC_DATA:
stmt = c_parser_oacc_data (loc, parser);
break;
case PRAGMA_OACC_HOST_DATA:
stmt = c_parser_oacc_host_data (loc, parser);
break;
case PRAGMA_OACC_KERNELS:
case PRAGMA_OACC_PARALLEL:
strcpy (p_name, "#pragma acc");

1
gcc/c/c-tree.h
View File

@ -653,6 +653,7 @@ extern tree c_finish_goto_ptr (location_t, tree);
extern tree c_expr_to_decl (tree, bool *, bool *);
extern tree c_finish_omp_construct (location_t, enum tree_code, tree, tree);
extern tree c_finish_oacc_data (location_t, tree, tree);
extern tree c_finish_oacc_host_data (location_t, tree, tree);
extern tree c_begin_omp_parallel (void);
extern tree c_finish_omp_parallel (location_t, tree, tree);
extern tree c_begin_omp_task (void);

20
gcc/c/c-typeck.c
View File

@ -11631,6 +11631,25 @@ c_finish_oacc_data (location_t loc, tree clauses, tree block)
return add_stmt (stmt);
}
/* Generate OACC_HOST_DATA, with CLAUSES and BLOCK as its compound
statement. LOC is the location of the OACC_HOST_DATA. */
tree
c_finish_oacc_host_data (location_t loc, tree clauses, tree block)
{
tree stmt;
block = c_end_compound_stmt (loc, block, true);
stmt = make_node (OACC_HOST_DATA);
TREE_TYPE (stmt) = void_type_node;
OACC_HOST_DATA_CLAUSES (stmt) = clauses;
OACC_HOST_DATA_BODY (stmt) = block;
SET_EXPR_LOCATION (stmt, loc);
return add_stmt (stmt);
}
/* Like c_begin_compound_stmt, except force the retention of the BLOCK. */
tree
@ -13074,6 +13093,7 @@ c_finish_omp_clauses (tree clauses, bool is_omp, bool declare_simd)
bitmap_set_bit (&map_head, DECL_UID (t));
goto check_dup_generic;
case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_IS_DEVICE_PTR:
case OMP_CLAUSE_USE_DEVICE_PTR:
t = OMP_CLAUSE_DECL (c);

14
gcc/cp/ChangeLog
View File

@ -1,3 +1,17 @@
2015-12-01 Julian Brown <julian@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
James Norris <James_Norris@mentor.com>
* cp-tree.h (finish_oacc_host_data): Add prototype.
* parser.c (cp_parser_omp_clause_name): Add use_device support.
(cp_parser_oacc_all_clauses): Add use_device support.
(OACC_HOST_DATA_CLAUSE_MASK): New macro.
(cp_parser_oacc_host_data): New function.
(cp_parser_omp_construct): Add host_data support.
(cp_parser_pragma): Add host_data support.
* semantics.c (finish_omp_clauses): Add use_device support.
(finish_oacc_host_data): New function.
2015-11-27 Martin Liska <mliska@suse.cz>
PR c++/68312

1
gcc/cp/cp-tree.h
View File

@ -6360,6 +6360,7 @@ extern void finish_omp_threadprivate (tree);
extern tree begin_omp_structured_block (void);
extern tree finish_omp_structured_block (tree);
extern tree finish_oacc_data (tree, tree);
extern tree finish_oacc_host_data (tree, tree);
extern tree finish_omp_construct (enum tree_code, tree, tree);
extern tree begin_omp_parallel (void);
extern tree finish_omp_parallel (tree, tree);

35
gcc/cp/parser.c
View File

@ -29232,6 +29232,8 @@ cp_parser_omp_clause_name (cp_parser *parser)
result = PRAGMA_OMP_CLAUSE_UNTIED;
else if (!strcmp ("use_device_ptr", p))
result = PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR;
else if (!strcmp ("use_device", p))
result = PRAGMA_OACC_CLAUSE_USE_DEVICE;
break;
case 'v':
if (!strcmp ("vector", p))
@ -31598,6 +31600,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 = "self";
break;
case PRAGMA_OACC_CLAUSE_USE_DEVICE:
clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_USE_DEVICE,
clauses);
c_name = "use_device";
break;
case PRAGMA_OACC_CLAUSE_SEQ:
clauses = cp_parser_oacc_simple_clause (parser, OMP_CLAUSE_SEQ,
clauses, here);
@ -34509,6 +34516,30 @@ cp_parser_oacc_data (cp_parser *parser, cp_token *pragma_tok)
return stmt;
}
#define OACC_HOST_DATA_CLAUSE_MASK \
( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_USE_DEVICE) )
/* OpenACC 2.0:
# pragma acc host_data <clauses> new-line
structured-block */
static tree
cp_parser_oacc_host_data (cp_parser *parser, cp_token *pragma_tok)
{
tree stmt, clauses, block;
unsigned int save;
clauses = cp_parser_oacc_all_clauses (parser, OACC_HOST_DATA_CLAUSE_MASK,
"#pragma acc host_data", pragma_tok);
block = begin_omp_parallel ();
save = cp_parser_begin_omp_structured_block (parser);
cp_parser_statement (parser, NULL_TREE, false, NULL);
cp_parser_end_omp_structured_block (parser, save);
stmt = finish_oacc_host_data (clauses, block);
return stmt;
}
/* OpenACC 2.0:
# pragma acc declare oacc-data-clause[optseq] new-line
*/
@ -36068,6 +36099,9 @@ cp_parser_omp_construct (cp_parser *parser, cp_token *pragma_tok)
case PRAGMA_OACC_EXIT_DATA:
stmt = cp_parser_oacc_enter_exit_data (parser, pragma_tok, false);
break;
case PRAGMA_OACC_HOST_DATA:
stmt = cp_parser_oacc_host_data (parser, pragma_tok);
break;
case PRAGMA_OACC_KERNELS:
case PRAGMA_OACC_PARALLEL:
strcpy (p_name, "#pragma acc");
@ -36645,6 +36679,7 @@ cp_parser_pragma (cp_parser *parser, enum pragma_context context)
case PRAGMA_OACC_DATA:
case PRAGMA_OACC_ENTER_DATA:
case PRAGMA_OACC_EXIT_DATA:
case PRAGMA_OACC_HOST_DATA:
case PRAGMA_OACC_KERNELS:
case PRAGMA_OACC_PARALLEL:
case PRAGMA_OACC_LOOP:

19
gcc/cp/semantics.c
View File

@ -6835,6 +6835,7 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd)
}
break;
case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_IS_DEVICE_PTR:
case OMP_CLAUSE_USE_DEVICE_PTR:
field_ok = allow_fields;
@ -7390,6 +7391,24 @@ finish_oacc_data (tree clauses, tree block)
return add_stmt (stmt);
}
/* Generate OACC_HOST_DATA, with CLAUSES and BLOCK as its compound
statement. */
tree
finish_oacc_host_data (tree clauses, tree block)
{
tree stmt;
block = finish_omp_structured_block (block);
stmt = make_node (OACC_HOST_DATA);
TREE_TYPE (stmt) = void_type_node;
OACC_HOST_DATA_CLAUSES (stmt) = clauses;
OACC_HOST_DATA_BODY (stmt) = block;
return add_stmt (stmt);
}
/* Generate OMP construct CODE, with BODY and CLAUSES as its compound
statement. */

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

@ -1356,6 +1356,9 @@ dump_gimple_omp_target (pretty_printer *buffer, gomp_target *gs,
case GF_OMP_TARGET_KIND_OACC_DECLARE:
kind = " oacc_declare";
break;
case GF_OMP_TARGET_KIND_OACC_HOST_DATA:
kind = " oacc_host_data";
break;
default:
gcc_unreachable ();
}

2
gcc/gimple.h
View File

@ -171,6 +171,7 @@ enum gf_mask {
GF_OMP_TARGET_KIND_OACC_UPDATE = 8,
GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA = 9,
GF_OMP_TARGET_KIND_OACC_DECLARE = 10,
GF_OMP_TARGET_KIND_OACC_HOST_DATA = 11,
/* 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_UPDATE:
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
case GF_OMP_TARGET_KIND_OACC_DECLARE:
case GF_OMP_TARGET_KIND_OACC_HOST_DATA:
return true;
default:
return false;

38
gcc/gimplify.c
View File

@ -122,6 +122,7 @@ enum omp_region_type
ORT_ACC_DATA = ORT_ACC | ORT_TARGET_DATA, /* Data construct. */
ORT_ACC_PARALLEL = ORT_ACC | ORT_TARGET, /* Parallel construct */
ORT_ACC_KERNELS = ORT_ACC | ORT_TARGET | 0x80, /* Kernels construct. */
ORT_ACC_HOST_DATA = ORT_ACC | ORT_TARGET_DATA | 0x80, /* Host data. */
/* Dummy OpenMP region, used to disable expansion of
DECL_VALUE_EXPRs in taskloop pre body. */
@ -6120,6 +6121,9 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
(splay_tree_key) decl);
if (n2)
{
if (octx->region_type == ORT_ACC_HOST_DATA)
error ("variable %qE declared in enclosing "
"%<host_data%> region", DECL_NAME (decl));
nflags |= GOVD_MAP;
goto found_outer;
}
@ -6418,6 +6422,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
case OMP_TARGET_DATA:
case OMP_TARGET_ENTER_DATA:
case OMP_TARGET_EXIT_DATA:
case OACC_HOST_DATA:
ctx->target_firstprivatize_array_bases = true;
default:
break;
@ -6683,6 +6688,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
case OMP_TARGET_DATA:
case OMP_TARGET_ENTER_DATA:
case OMP_TARGET_EXIT_DATA:
case OACC_HOST_DATA:
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
|| (OMP_CLAUSE_MAP_KIND (c)
== GOMP_MAP_FIRSTPRIVATE_REFERENCE))
@ -6695,6 +6701,22 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
}
if (remove)
break;
if (DECL_P (decl) && outer_ctx && (region_type & ORT_ACC))
{
struct gimplify_omp_ctx *octx;
for (octx = outer_ctx; octx; octx = octx->outer_context)
{
if (octx->region_type != ORT_ACC_HOST_DATA)
break;
splay_tree_node n2
= splay_tree_lookup (octx->variables,
(splay_tree_key) decl);
if (n2)
error_at (OMP_CLAUSE_LOCATION (c), "variable %qE "
"declared in enclosing %<host_data%> region",
DECL_NAME (decl));
}
}
if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
OMP_CLAUSE_SIZE (c) = DECL_P (decl) ? DECL_SIZE_UNIT (decl)
: TYPE_SIZE_UNIT (TREE_TYPE (decl));
@ -7092,6 +7114,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
}
goto do_notice;
case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_USE_DEVICE_PTR:
flags = GOVD_FIRSTPRIVATE | GOVD_EXPLICIT;
goto do_add;
@ -7327,7 +7350,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
break;
case OMP_CLAUSE_DEVICE_RESIDENT:
case OMP_CLAUSE_USE_DEVICE:
remove = true;
break;
@ -9365,6 +9387,9 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
case OMP_TEAMS:
ort = OMP_TEAMS_COMBINED (expr) ? ORT_COMBINED_TEAMS : ORT_TEAMS;
break;
case OACC_HOST_DATA:
ort = ORT_ACC_HOST_DATA;
break;
default:
gcc_unreachable ();
}
@ -9386,6 +9411,7 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
switch (TREE_CODE (expr))
{
case OACC_DATA:
case OACC_HOST_DATA:
end_ix = BUILT_IN_GOACC_DATA_END;
break;
case OMP_TARGET_DATA:
@ -9418,6 +9444,10 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_KERNELS,
OMP_CLAUSES (expr));
break;
case OACC_HOST_DATA:
stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_HOST_DATA,
OMP_CLAUSES (expr));
break;
case OACC_PARALLEL:
stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_PARALLEL,
OMP_CLAUSES (expr));
@ -10527,16 +10557,12 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
ret = GS_ALL_DONE;
break;
case OACC_HOST_DATA:
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_HOST_DATA:
case OACC_DATA:
case OACC_KERNELS:
case OACC_PARALLEL:

2
gcc/omp-builtins.def
View File

@ -47,6 +47,8 @@ DEF_GOACC_BUILTIN (BUILT_IN_GOACC_UPDATE, "GOACC_update",
DEF_GOACC_BUILTIN (BUILT_IN_GOACC_WAIT, "GOACC_wait",
BT_FN_VOID_INT_INT_VAR,
ATTR_NOTHROW_LIST)
DEF_GOACC_BUILTIN (BUILT_IN_GOACC_HOST_DATA, "GOACC_host_data",
BT_FN_VOID_INT_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
DEF_GOACC_BUILTIN_COMPILER (BUILT_IN_ACC_ON_DEVICE, "acc_on_device",
BT_FN_INT_INT, ATTR_CONST_NOTHROW_LEAF_LIST)

25
gcc/omp-low.c
View File

@ -1942,6 +1942,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
install_var_local (decl, ctx);
break;
case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_USE_DEVICE_PTR:
decl = OMP_CLAUSE_DECL (c);
if (TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
@ -2144,7 +2145,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
break;
case OMP_CLAUSE_DEVICE_RESIDENT:
case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE__CACHE_:
sorry ("Clause not supported yet");
break;
@ -2295,6 +2295,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_SIMD:
case OMP_CLAUSE_NOGROUP:
case OMP_CLAUSE_DEFAULTMAP:
case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE__CILK_FOR_COUNT_:
case OMP_CLAUSE_ASYNC:
@ -2312,7 +2313,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
break;
case OMP_CLAUSE_DEVICE_RESIDENT:
case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE__CACHE_:
sorry ("Clause not supported yet");
break;
@ -3615,6 +3615,8 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx)
case GF_OMP_TARGET_KIND_OACC_UPDATE: stmt_name = "update"; break;
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
stmt_name = "enter/exit data"; break;
case GF_OMP_TARGET_KIND_OACC_HOST_DATA: stmt_name = "host_data";
break;
default: gcc_unreachable ();
}
switch (gimple_omp_target_kind (ctx->stmt))
@ -3626,6 +3628,8 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx)
case GF_OMP_TARGET_KIND_OACC_KERNELS:
ctx_stmt_name = "kernels"; break;
case GF_OMP_TARGET_KIND_OACC_DATA: ctx_stmt_name = "data"; break;
case GF_OMP_TARGET_KIND_OACC_HOST_DATA:
ctx_stmt_name = "host_data"; break;
default: gcc_unreachable ();
}
@ -12508,6 +12512,7 @@ expand_omp_target (struct omp_region *region)
break;
case GF_OMP_TARGET_KIND_DATA:
case GF_OMP_TARGET_KIND_OACC_DATA:
case GF_OMP_TARGET_KIND_OACC_HOST_DATA:
data_region = true;
break;
default:
@ -12751,6 +12756,9 @@ expand_omp_target (struct omp_region *region)
case GF_OMP_TARGET_KIND_OACC_DECLARE:
start_ix = BUILT_IN_GOACC_DECLARE;
break;
case GF_OMP_TARGET_KIND_OACC_HOST_DATA:
start_ix = BUILT_IN_GOACC_HOST_DATA;
break;
default:
gcc_unreachable ();
}
@ -12875,6 +12883,7 @@ expand_omp_target (struct omp_region *region)
case BUILT_IN_GOACC_DATA_START:
case BUILT_IN_GOACC_DECLARE:
case BUILT_IN_GOMP_TARGET_DATA:
case BUILT_IN_GOACC_HOST_DATA:
break;
case BUILT_IN_GOMP_TARGET:
case BUILT_IN_GOMP_TARGET_UPDATE:
@ -13182,6 +13191,7 @@ build_omp_regions_1 (basic_block bb, struct omp_region *parent,
case GF_OMP_TARGET_KIND_OACC_PARALLEL:
case GF_OMP_TARGET_KIND_OACC_KERNELS:
case GF_OMP_TARGET_KIND_OACC_DATA:
case GF_OMP_TARGET_KIND_OACC_HOST_DATA:
break;
case GF_OMP_TARGET_KIND_UPDATE:
case GF_OMP_TARGET_KIND_ENTER_DATA:
@ -14982,6 +14992,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
break;
case GF_OMP_TARGET_KIND_DATA:
case GF_OMP_TARGET_KIND_OACC_DATA:
case GF_OMP_TARGET_KIND_OACC_HOST_DATA:
data_region = true;
break;
default:
@ -15188,6 +15199,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
}
break;
case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_IS_DEVICE_PTR:
var = OMP_CLAUSE_DECL (c);
@ -15573,12 +15585,14 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
build_int_cstu (tkind_type, tkind));
break;
case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_IS_DEVICE_PTR:
ovar = OMP_CLAUSE_DECL (c);
var = lookup_decl_in_outer_ctx (ovar, ctx);
x = build_sender_ref (ovar, ctx);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR)
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE)
tkind = GOMP_MAP_USE_DEVICE_PTR;
else
tkind = GOMP_MAP_FIRSTPRIVATE_INT;
@ -15781,10 +15795,12 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
gimple_build_assign (new_var, x));
}
break;
case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_IS_DEVICE_PTR:
var = OMP_CLAUSE_DECL (c);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR)
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE)
x = build_sender_ref (var, ctx);
else
x = build_receiver_ref (var, false, ctx);
@ -16771,6 +16787,7 @@ make_gimple_omp_edges (basic_block bb, struct omp_region **region,
case GF_OMP_TARGET_KIND_OACC_PARALLEL:
case GF_OMP_TARGET_KIND_OACC_KERNELS:
case GF_OMP_TARGET_KIND_OACC_DATA:
case GF_OMP_TARGET_KIND_OACC_HOST_DATA:
break;
case GF_OMP_TARGET_KIND_UPDATE:
case GF_OMP_TARGET_KIND_ENTER_DATA:

2
gcc/tree-nested.c
View File

@ -1072,6 +1072,7 @@ convert_nonlocal_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
case OMP_CLAUSE_SHARED:
case OMP_CLAUSE_TO_DECLARE:
case OMP_CLAUSE_LINK:
case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_IS_DEVICE_PTR:
do_decl_clause:
@ -1743,6 +1744,7 @@ convert_local_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
case OMP_CLAUSE_SHARED:
case OMP_CLAUSE_TO_DECLARE:
case OMP_CLAUSE_LINK:
case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_IS_DEVICE_PTR:
do_decl_clause:

12
libgomp/ChangeLog
View File

@ -1,3 +1,15 @@
2015-12-01 Julian Brown <julian@codesourcery.com>
James Norris <James_Norris@mentor.com>
* oacc-parallel.c (GOACC_host_data): New function.
* libgomp.map (GOACC_host_data): Add to GOACC_2.0.1.
* testsuite/libgomp.oacc-c-c++-common/host_data-1.c: New test.
* testsuite/libgomp.oacc-c-c++-common/host_data-2.c: New test.
* testsuite/libgomp.oacc-c-c++-common/host_data-3.c: New test.
* testsuite/libgomp.oacc-c-c++-common/host_data-4.c: New test.
* testsuite/libgomp.oacc-c-c++-common/host_data-5.c: New test.
* testsuite/libgomp.oacc-c-c++-common/host_data-6.c: New test.
2015-11-30 James Norris <jnorris@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>

1
libgomp/libgomp.map
View File

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

40
libgomp/oacc-parallel.c
View File

@ -490,6 +490,46 @@ GOACC_wait (int async, int num_waits, ...)
goacc_thread ()->dev->openacc.async_wait_all_async_func (acc_async_noval);
}
void
GOACC_host_data (int device, size_t mapnum,
void **hostaddrs, size_t *sizes, unsigned short *kinds)
{
bool host_fallback = device == GOMP_DEVICE_HOST_FALLBACK;
struct target_mem_desc *tgt;
#ifdef HAVE_INTTYPES_H
gomp_debug (0, "%s: mapnum=%"PRIu64", hostaddrs=%p, size=%p, kinds=%p\n",
__FUNCTION__, (uint64_t) mapnum, hostaddrs, sizes, kinds);
#else
gomp_debug (0, "%s: mapnum=%lu, hostaddrs=%p, sizes=%p, kinds=%p\n",
__FUNCTION__, (unsigned long) mapnum, hostaddrs, sizes, kinds);
#endif
goacc_lazy_initialize ();
struct goacc_thread *thr = goacc_thread ();
struct gomp_device_descr *acc_dev = thr->dev;
/* Host fallback or 'do nothing'. */
if ((acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
|| host_fallback)
{
tgt = gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, true,
GOMP_MAP_VARS_OPENACC);
tgt->prev = thr->mapped_data;
thr->mapped_data = tgt;
return;
}
gomp_debug (0, " %s: prepare mappings\n", __FUNCTION__);
tgt = gomp_map_vars (acc_dev, mapnum, hostaddrs, NULL, sizes, kinds, true,
GOMP_MAP_VARS_OPENACC);
gomp_debug (0, " %s: mappings prepared\n", __FUNCTION__);
tgt->prev = thr->mapped_data;
thr->mapped_data = tgt;
}
int
GOACC_get_num_threads (void)
{

100
libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c Normal file
View File

@ -0,0 +1,100 @@
/* { dg-do run { target openacc_nvidia_accel_selected } } */
/* { dg-additional-options "-lcuda -lcublas -lcudart" } */
#include <stdlib.h>
#include <openacc.h>
#include <cuda.h>
#include <cuda_runtime_api.h>
#include <cublas_v2.h>
void
saxpy_host (int n, float a, float *x, float *y)
{
int i;
for (i = 0; i < n; i++)
y[i] = y[i] + a * x[i];
}
#pragma acc routine
void
saxpy_target (int n, float a, float *x, float *y)
{
int i;
for (i = 0; i < n; i++)
y[i] = y[i] + a * x[i];
}
int
main(int argc, char **argv)
{
#define N 8
int i;
float x_ref[N], y_ref[N];
float x[N], y[N];
cublasHandle_t h;
float a = 2.0;
for (i = 0; i < N; i++)
{
x[i] = x_ref[i] = 4.0 + i;
y[i] = y_ref[i] = 3.0;
}
saxpy_host (N, a, x_ref, y_ref);
cublasCreate (&h);
#pragma acc data copyin (x[0:N]) copy (y[0:N])
{
#pragma acc host_data use_device (x, y)
{
cublasSaxpy (h, N, &a, x, 1, y, 1);
}
}
for (i = 0; i < N; i++)
{
if (y[i] != y_ref[i])
abort ();
}
#pragma acc data create (x[0:N]) copyout (y[0:N])
{
#pragma acc kernels
for (i = 0; i < N; i++)
y[i] = 3.0;
#pragma acc host_data use_device (x, y)
{
cublasSaxpy (h, N, &a, x, 1, y, 1);
}
}
cublasDestroy (h);
for (i = 0; i < N; i++)
{
if (y[i] != y_ref[i])
abort ();
}
for (i = 0; i < N; i++)
y[i] = 3.0;
/* There's no need to use host_data here. */
#pragma acc data copyin (x[0:N]) copyin (a) copy (y[0:N])
{
#pragma acc parallel present (x[0:N]) pcopy (y[0:N]) present (a)
saxpy_target (N, a, x, y);
}
for (i = 0; i < N; i++)
{
if (y[i] != y_ref[i])
abort ();
}
return 0;
}

31
libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-2.c Normal file
View File

@ -0,0 +1,31 @@
/* { dg-do run } */
#include <stdlib.h>
#include <openacc.h>
char *global_in_host;
void foo (char *in)
{
if (!acc_is_present (global_in_host, sizeof (*global_in_host))
|| in != acc_deviceptr (global_in_host))
abort ();
}
int
main (int argc, char **argv)
{
char mydata[1024];
global_in_host = mydata;
#pragma acc data copyin(mydata)
{
#pragma acc host_data use_device (mydata)
{
foo (mydata);
}
}
return 0;
}

29
libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-3.c Normal file
View File

@ -0,0 +1,29 @@
/* { dg-do compile } */
#include <openacc.h>
#include <stdlib.h>
#define N 1024
int main (int argc, char* argv[])
{
int x[N];
#pragma acc data copyin (x[0:N])
{
int *xp;
#pragma acc host_data use_device (x)
{
/* This use of the present clause is undefined behaviour for OpenACC. */
#pragma acc parallel present (x) copyout (xp) /* { dg-error "variable 'x' declared in enclosing 'host_data' region" } */
{
xp = x;
}
}
if (xp != acc_deviceptr (x))
abort ();
}
return 0;
}

29
libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-4.c Normal file
View File

@ -0,0 +1,29 @@
/* { dg-do run } */
#include <openacc.h>
#include <stdlib.h>
#define N 1024
int main (int argc, char* argv[])
{
int x[N], *xp2;
#pragma acc data copyin (x[0:N])
{
int *xp;
#pragma acc host_data use_device (x)
{
#pragma acc data
{
xp = x;
}
xp2 = x;
}
if (xp != acc_deviceptr (x) || xp2 != xp)
abort ();
}
return 0;
}

38
libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-5.c Normal file
View File

@ -0,0 +1,38 @@
/* { dg-do run } */
#include <openacc.h>
#include <stdlib.h>
#define N 1024
int main (int argc, char* argv[])
{
int x[N], y[N], *yp;
yp = y + 1;
#pragma acc data copyin (x[0:N])
{
int *xp, *yp2;
#pragma acc host_data use_device (x)
{
#pragma acc data copyin (y)
{
#pragma acc host_data use_device (yp)
{
xp = x;
yp2 = yp;
}
if (yp2 != acc_deviceptr (yp))
abort ();
}
}
if (xp != acc_deviceptr (x))
abort ();
}
return 0;
}

31
libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-6.c Normal file
View File

@ -0,0 +1,31 @@
/* { dg-do compile } */
#include <openacc.h>
#include <stdlib.h>
#define N 1024
int main (int argc, char* argv[])
{
int x[N];
#pragma acc data copyin (x[0:N])
{
int *xp;
#pragma acc host_data use_device (x)
{
/* Here 'x' being implicitly firstprivate for the parallel region
conflicts with it being declared as use_device in the enclosing
host_data region. */
#pragma acc parallel copyout (xp)
{
xp = x; /* { dg-error "variable 'x' declared in enclosing 'host_data' region" } */
}
}
if (xp != acc_deviceptr (x))
abort ();
}
return 0;
}