commit ac77cadbe27ca43d036f943c48fb4bf59bc84b36
Author: Julian Brown <julian@codesourcery.com>
Date: Mon Nov 2 06:31:47 2015 -0800
OpenACC host_data support using mapping regions.
@@ -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 },
@@ -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,
@@ -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);
@@ -13590,6 +13605,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
structured-block
@@ -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");
@@ -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);
@@ -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);
@@ -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);
@@ -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:
@@ -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. */
@@ -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 ();
}
@@ -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;
@@ -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_TARGET_DATA | ORT_TARGET)))
+ break;
+ splay_tree_node n2
+ = splay_tree_lookup (octx->variables,
+ (splay_tree_key) decl);
+ if (n2 && octx->region_type == ORT_ACC_HOST_DATA)
+ 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:
@@ -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)
@@ -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:
@@ -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:
@@ -1721,6 +1722,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:
@@ -394,6 +394,7 @@ GOACC_2.0.1 {
global:
GOACC_declare;
GOACC_parallel_keyed;
+ GOACC_host_data;
} GOACC_2.0;
GOMP_PLUGIN_1.0 {
@@ -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)
{
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}