mirror of git://gcc.gnu.org/git/gcc.git
OpenMP: C++ "declare mapper" support
This patch adds support for OpenMP 5.0 "declare mapper" functionality
for C++. I've merged it to og13 based on the last version
posted upstream, with some minor changes due to the newly-added
'present' map modifier support. There's also a fix to splay-tree
traversal in gimplify.cc:omp_instantiate_implicit_mappers, and this patch
omits the rearrangement of gimplify.cc:gimplify_{scan,adjust}_omp_clauses
that I separated out into its own patch and applied (to og13) already.
gcc/c-family/
* c-common.h (c_omp_region_type): Add C_ORT_DECLARE_MAPPER and
C_ORT_OMP_DECLARE_MAPPER codes.
(omp_mapper_list): Add forward declaration.
(c_omp_find_nested_mappers, c_omp_instantiate_mappers): Add prototypes.
* c-omp.cc (c_omp_find_nested_mappers): New function.
(remap_mapper_decl_info): New struct.
(remap_mapper_decl_1, omp_instantiate_mapper,
c_omp_instantiate_mappers): New functions.
gcc/cp/
* constexpr.cc (reduced_constant_expression_p): Add OMP_DECLARE_MAPPER
case.
(cxx_eval_constant_expression, potential_constant_expression_1):
Likewise.
* cp-gimplify.cc (cxx_omp_finish_mapper_clauses): New function.
* cp-objcp-common.h (LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES,
LANG_HOOKS_OMP_MAPPER_LOOKUP, LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE,
LANG_HOOKS_OMP_MAP_ARRAY_SECTION): Define langhooks.
* cp-tree.h (lang_decl_base): Add omp_declare_mapper_p field. Recount
spare bits comment.
(DECL_OMP_DECLARE_MAPPER_P): New macro.
(omp_mapper_id): Add prototype.
(cp_check_omp_declare_mapper): Add prototype.
(omp_instantiate_mappers): Add prototype.
(cxx_omp_finish_mapper_clauses): Add prototype.
(cxx_omp_mapper_lookup): Add prototype.
(cxx_omp_extract_mapper_directive): Add prototype.
(cxx_omp_map_array_section): Add prototype.
* decl.cc (check_initializer): Add OpenMP declare mapper support.
(cp_finish_decl): Set DECL_INITIAL for OpenMP declare mapper var decls
as appropriate.
* decl2.cc (mark_used): Instantiate OpenMP "declare mapper" magic var
decls.
* error.cc (dump_omp_declare_mapper): New function.
(dump_simple_decl): Use above.
* parser.cc (cp_parser_omp_clause_map): Add KIND parameter. Support
"mapper" modifier.
(cp_parser_omp_all_clauses): Add KIND argument to
cp_parser_omp_clause_map call.
(cp_parser_omp_target): Call omp_instantiate_mappers before
finish_omp_clauses.
(cp_parser_omp_declare_mapper): New function.
(cp_parser_omp_declare): Add "declare mapper" support.
* pt.cc (tsubst_decl): Adjust name of "declare mapper" magic var decls
once we know their type.
(tsubst_omp_clauses): Call omp_instantiate_mappers before
finish_omp_clauses, for target regions.
(tsubst_expr): Support OMP_DECLARE_MAPPER nodes.
(instantiate_decl): Instantiate initialiser (i.e definition) for OpenMP
declare mappers.
* semantics.cc (gimplify.h): Include.
(omp_mapper_id, omp_mapper_lookup, omp_extract_mapper_directive,
cxx_omp_map_array_section, cp_check_omp_declare_mapper): New functions.
(finish_omp_clauses): Delete GOMP_MAP_PUSH_MAPPER_NAME and
GOMP_MAP_POP_MAPPER_NAME artificial clauses.
(omp_target_walk_data): Add MAPPERS field.
(finish_omp_target_clauses_r): Scan for uses of struct/union/class type
variables.
(finish_omp_target_clauses): Create artificial mapper binding clauses
for used structs/unions/classes in offload region.
gcc/fortran/
* parse.cc (tree.h, fold-const.h, tree-hash-traits.h): Add includes
(for additions to omp-general.h).
gcc/
* gimplify.cc (gimplify_omp_ctx): Add IMPLICIT_MAPPERS field.
(new_omp_context): Initialise IMPLICIT_MAPPERS hash map.
(delete_omp_context): Delete IMPLICIT_MAPPERS hash map.
(instantiate_mapper_info): New structs.
(remap_mapper_decl_1, omp_mapper_copy_decl, omp_instantiate_mapper,
omp_instantiate_implicit_mappers): New functions.
(gimplify_scan_omp_clauses): Handle MAPPER_BINDING clauses.
(gimplify_adjust_omp_clauses): Instantiate implicit declared mappers.
(gimplify_omp_declare_mapper): New function.
(gimplify_expr): Call above function.
* langhooks-def.h (lhd_omp_mapper_lookup,
lhd_omp_extract_mapper_directive, lhd_omp_map_array_section): Add
prototypes.
(LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES,
LANG_HOOKS_OMP_MAPPER_LOOKUP, LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE,
LANG_HOOKS_OMP_MAP_ARRAY_SECTION): Define macros.
(LANG_HOOK_DECLS): Add above macros.
* langhooks.cc (lhd_omp_mapper_lookup,
lhd_omp_extract_mapper_directive, lhd_omp_map_array_section): New
dummy functions.
* langhooks.h (lang_hooks_for_decls): Add OMP_FINISH_MAPPER_CLAUSES,
OMP_MAPPER_LOOKUP, OMP_EXTRACT_MAPPER_DIRECTIVE, OMP_MAP_ARRAY_SECTION
hooks.
* omp-general.h (omp_name_type<T>): Add templatized struct, hash type
traits (for omp_name_type<tree> specialization).
(omp_mapper_list<T>): Add struct.
* tree-core.h (omp_clause_code): Add OMP_CLAUSE__MAPPER_BINDING_.
* tree-pretty-print.cc (dump_omp_clause): Support GOMP_MAP_UNSET,
GOMP_MAP_PUSH_MAPPER_NAME, GOMP_MAP_POP_MAPPER_NAME artificial mapping
clauses. Support OMP_CLAUSE__MAPPER_BINDING_ and OMP_DECLARE_MAPPER.
* tree.cc (omp_clause_num_ops, omp_clause_code_name): Add
OMP_CLAUSE__MAPPER_BINDING_.
* tree.def (OMP_DECLARE_MAPPER): New tree code.
* tree.h (OMP_DECLARE_MAPPER_ID, OMP_DECLARE_MAPPER_DECL,
OMP_DECLARE_MAPPER_CLAUSES): New defines.
(OMP_CLAUSE__MAPPER_BINDING__ID, OMP_CLAUSE__MAPPER_BINDING__DECL,
OMP_CLAUSE__MAPPER_BINDING__MAPPER): New defines.
include/
* gomp-constants.h (gomp_map_kind): Add GOMP_MAP_UNSET,
GOMP_MAP_PUSH_MAPPER_NAME, GOMP_MAP_POP_MAPPER_NAME artificial mapping
clause types.
gcc/testsuite/
* c-c++-common/gomp/map-6.c: Update error scan output.
* c-c++-common/gomp/declare-mapper-3.c: New test (only enabled for C++
for now).
* c-c++-common/gomp/declare-mapper-4.c: Likewise.
* c-c++-common/gomp/declare-mapper-5.c: Likewise.
* c-c++-common/gomp/declare-mapper-6.c: Likewise.
* c-c++-common/gomp/declare-mapper-7.c: Likewise.
* c-c++-common/gomp/declare-mapper-8.c: Likewise.
* c-c++-common/gomp/declare-mapper-9.c: Likewise.
* c-c++-common/gomp/declare-mapper-10.c: Likewise.
* c-c++-common/gomp/declare-mapper-12.c: Likewise.
* g++.dg/gomp/declare-mapper-1.C: New test.
* g++.dg/gomp/declare-mapper-2.C: New test.
* g++.dg/gomp/declare-mapper-3.C: New test.
libgomp/
* testsuite/libgomp.c++/declare-mapper-1.C: New test.
* testsuite/libgomp.c++/declare-mapper-2.C: New test.
* testsuite/libgomp.c++/declare-mapper-3.C: New test.
* testsuite/libgomp.c++/declare-mapper-4.C: New test.
* testsuite/libgomp.c++/declare-mapper-5.C: New test.
* testsuite/libgomp.c++/declare-mapper-6.C: New test.
* testsuite/libgomp.c++/declare-mapper-7.C: New test.
* testsuite/libgomp.c++/declare-mapper-8.C: New test.
* testsuite/libgomp.c-c++-common/declare-mapper-9.c: New test (only
enabled for C++ for now).
* testsuite/libgomp.c-c++-common/declare-mapper-10.c: Likewise.
* testsuite/libgomp.c-c++-common/declare-mapper-11.c: Likewise.
* testsuite/libgomp.c-c++-common/declare-mapper-12.c: Likewise.
* testsuite/libgomp.c-c++-common/declare-mapper-13.c: Likewise.
* testsuite/libgomp.c-c++-common/declare-mapper-14.c: Likewise.
Co-authored-by: Tobias Burnus <tburnus@baylibre.com>
This commit is contained in:
Julian Brown
Tobias Burnus
parent
0e0f963bcf
commit
48973e8783
|
|
@ -1303,10 +1303,12 @@ enum c_omp_region_type
|
|||
C_ORT_TARGET = 1 << 3,
|
||||
C_ORT_EXIT_DATA = 1 << 4,
|
||||
C_ORT_INTEROP = 1 << 5,
|
||||
C_ORT_DECLARE_MAPPER = 1 << 6,
|
||||
C_ORT_OMP_DECLARE_SIMD = C_ORT_OMP | C_ORT_DECLARE_SIMD,
|
||||
C_ORT_OMP_TARGET = C_ORT_OMP | C_ORT_TARGET,
|
||||
C_ORT_OMP_EXIT_DATA = C_ORT_OMP | C_ORT_EXIT_DATA,
|
||||
C_ORT_OMP_INTEROP = C_ORT_OMP | C_ORT_INTEROP,
|
||||
C_ORT_OMP_DECLARE_MAPPER = C_ORT_OMP | C_ORT_DECLARE_MAPPER,
|
||||
C_ORT_ACC_TARGET = C_ORT_ACC | C_ORT_TARGET
|
||||
};
|
||||
|
||||
|
|
@ -1345,6 +1347,9 @@ extern enum omp_clause_defaultmap_kind c_omp_predetermined_mapping (tree);
|
|||
extern tree c_omp_check_context_selector (location_t, tree);
|
||||
extern void c_omp_mark_declare_variant (location_t, tree, tree);
|
||||
extern void c_omp_adjust_map_clauses (tree, bool);
|
||||
template<typename T> struct omp_mapper_list;
|
||||
extern void c_omp_find_nested_mappers (struct omp_mapper_list<tree> *, tree);
|
||||
extern tree c_omp_instantiate_mappers (tree);
|
||||
|
||||
namespace omp_addr_tokenizer { struct omp_addr_token; }
|
||||
typedef omp_addr_tokenizer::omp_addr_token omp_addr_token;
|
||||
|
|
|
|||
|
|
@ -4282,6 +4282,306 @@ c_omp_address_inspector::expand_map_clause (tree c, tree expr,
|
|||
return error_mark_node;
|
||||
}
|
||||
|
||||
/* Given a mapper function MAPPER_FN, recursively scan through the map clauses
|
||||
for that mapper, and if any of those should use a (named or unnamed) mapper
|
||||
themselves, add it to MLIST. */
|
||||
|
||||
void
|
||||
c_omp_find_nested_mappers (omp_mapper_list<tree> *mlist, tree mapper_fn)
|
||||
{
|
||||
tree mapper = lang_hooks.decls.omp_extract_mapper_directive (mapper_fn);
|
||||
tree mapper_name = NULL_TREE;
|
||||
|
||||
if (mapper == error_mark_node)
|
||||
return;
|
||||
|
||||
gcc_assert (TREE_CODE (mapper) == OMP_DECLARE_MAPPER);
|
||||
|
||||
for (tree clause = OMP_DECLARE_MAPPER_CLAUSES (mapper);
|
||||
clause;
|
||||
clause = OMP_CLAUSE_CHAIN (clause))
|
||||
{
|
||||
tree expr = OMP_CLAUSE_DECL (clause);
|
||||
enum gomp_map_kind clause_kind = OMP_CLAUSE_MAP_KIND (clause);
|
||||
tree elem_type;
|
||||
|
||||
if (clause_kind == GOMP_MAP_PUSH_MAPPER_NAME)
|
||||
{
|
||||
mapper_name = expr;
|
||||
continue;
|
||||
}
|
||||
else if (clause_kind == GOMP_MAP_POP_MAPPER_NAME)
|
||||
{
|
||||
mapper_name = NULL_TREE;
|
||||
continue;
|
||||
}
|
||||
|
||||
gcc_assert (TREE_CODE (expr) != TREE_LIST);
|
||||
if (TREE_CODE (expr) == OMP_ARRAY_SECTION)
|
||||
{
|
||||
while (TREE_CODE (expr) == OMP_ARRAY_SECTION)
|
||||
expr = TREE_OPERAND (expr, 0);
|
||||
|
||||
elem_type = TREE_TYPE (expr);
|
||||
}
|
||||
else
|
||||
elem_type = TREE_TYPE (expr);
|
||||
|
||||
/* This might be too much... or not enough? */
|
||||
while (TREE_CODE (elem_type) == ARRAY_TYPE
|
||||
|| TREE_CODE (elem_type) == POINTER_TYPE
|
||||
|| TREE_CODE (elem_type) == REFERENCE_TYPE)
|
||||
elem_type = TREE_TYPE (elem_type);
|
||||
|
||||
elem_type = TYPE_MAIN_VARIANT (elem_type);
|
||||
|
||||
if (RECORD_OR_UNION_TYPE_P (elem_type)
|
||||
&& !mlist->contains (mapper_name, elem_type))
|
||||
{
|
||||
tree nested_mapper_fn
|
||||
= lang_hooks.decls.omp_mapper_lookup (mapper_name, elem_type);
|
||||
|
||||
if (nested_mapper_fn)
|
||||
{
|
||||
mlist->add_mapper (mapper_name, elem_type, nested_mapper_fn);
|
||||
c_omp_find_nested_mappers (mlist, nested_mapper_fn);
|
||||
}
|
||||
else if (mapper_name)
|
||||
{
|
||||
error ("mapper %qE not found for type %qT", mapper_name,
|
||||
elem_type);
|
||||
continue;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
struct remap_mapper_decl_info
|
||||
{
|
||||
tree dummy_var;
|
||||
tree expr;
|
||||
};
|
||||
|
||||
/* Helper for rewriting DUMMY_VAR into EXPR in a map clause decl. */
|
||||
|
||||
static tree
|
||||
remap_mapper_decl_1 (tree *tp, int *walk_subtrees, void *data)
|
||||
{
|
||||
remap_mapper_decl_info *map_info = (remap_mapper_decl_info *) data;
|
||||
|
||||
if (operand_equal_p (*tp, map_info->dummy_var))
|
||||
{
|
||||
*tp = map_info->expr;
|
||||
*walk_subtrees = 0;
|
||||
}
|
||||
|
||||
return NULL_TREE;
|
||||
}
|
||||
|
||||
/* Instantiate a mapper MAPPER for expression EXPR, adding new clauses to
|
||||
OUTLIST. OUTER_KIND is the mapping kind to use if not already specified in
|
||||
the mapper declaration. */
|
||||
|
||||
static tree *
|
||||
omp_instantiate_mapper (tree *outlist, tree mapper, tree expr,
|
||||
enum gomp_map_kind outer_kind)
|
||||
{
|
||||
tree clauses = OMP_DECLARE_MAPPER_CLAUSES (mapper);
|
||||
tree dummy_var = OMP_DECLARE_MAPPER_DECL (mapper);
|
||||
tree mapper_name = NULL_TREE;
|
||||
|
||||
remap_mapper_decl_info map_info;
|
||||
map_info.dummy_var = dummy_var;
|
||||
map_info.expr = expr;
|
||||
|
||||
for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
|
||||
{
|
||||
tree unshared = unshare_expr (c);
|
||||
enum gomp_map_kind clause_kind = OMP_CLAUSE_MAP_KIND (c);
|
||||
tree t = OMP_CLAUSE_DECL (unshared);
|
||||
tree type = NULL_TREE;
|
||||
bool nonunit_array_with_mapper = false;
|
||||
|
||||
if (clause_kind == GOMP_MAP_PUSH_MAPPER_NAME)
|
||||
{
|
||||
mapper_name = t;
|
||||
continue;
|
||||
}
|
||||
else if (clause_kind == GOMP_MAP_POP_MAPPER_NAME)
|
||||
{
|
||||
mapper_name = NULL_TREE;
|
||||
continue;
|
||||
}
|
||||
|
||||
if (TREE_CODE (t) == OMP_ARRAY_SECTION)
|
||||
{
|
||||
location_t loc = OMP_CLAUSE_LOCATION (c);
|
||||
tree t2 = lang_hooks.decls.omp_map_array_section (loc, t);
|
||||
|
||||
if (t2 == t)
|
||||
{
|
||||
nonunit_array_with_mapper = true;
|
||||
/* We'd want use the mapper for the element type if this worked:
|
||||
look that one up. */
|
||||
type = TREE_TYPE (TREE_TYPE (t));
|
||||
}
|
||||
else
|
||||
{
|
||||
t = t2;
|
||||
type = TREE_TYPE (t);
|
||||
}
|
||||
}
|
||||
else
|
||||
type = TREE_TYPE (t);
|
||||
|
||||
gcc_assert (type);
|
||||
|
||||
if (type == error_mark_node)
|
||||
continue;
|
||||
|
||||
walk_tree (&unshared, remap_mapper_decl_1, &map_info, NULL);
|
||||
|
||||
if (OMP_CLAUSE_MAP_KIND (unshared) == GOMP_MAP_UNSET)
|
||||
OMP_CLAUSE_SET_MAP_KIND (unshared, outer_kind);
|
||||
|
||||
type = TYPE_MAIN_VARIANT (type);
|
||||
|
||||
tree mapper_fn = lang_hooks.decls.omp_mapper_lookup (mapper_name, type);
|
||||
|
||||
if (mapper_fn && nonunit_array_with_mapper)
|
||||
{
|
||||
sorry ("user-defined mapper with non-unit length array section");
|
||||
continue;
|
||||
}
|
||||
else if (mapper_fn)
|
||||
{
|
||||
tree nested_mapper
|
||||
= lang_hooks.decls.omp_extract_mapper_directive (mapper_fn);
|
||||
if (nested_mapper != mapper)
|
||||
{
|
||||
if (clause_kind == GOMP_MAP_UNSET)
|
||||
clause_kind = outer_kind;
|
||||
|
||||
outlist = omp_instantiate_mapper (outlist, nested_mapper,
|
||||
t, clause_kind);
|
||||
continue;
|
||||
}
|
||||
}
|
||||
else if (mapper_name)
|
||||
{
|
||||
error ("mapper %qE not found for type %qT", mapper_name, type);
|
||||
continue;
|
||||
}
|
||||
|
||||
*outlist = unshared;
|
||||
outlist = &OMP_CLAUSE_CHAIN (unshared);
|
||||
}
|
||||
|
||||
return outlist;
|
||||
}
|
||||
|
||||
/* Given a list of CLAUSES, scan each clause and invoke a user-defined mapper
|
||||
appropriate to the type of the data in that clause, if such a mapper is
|
||||
visible in the current parsing context. */
|
||||
|
||||
tree
|
||||
c_omp_instantiate_mappers (tree clauses)
|
||||
{
|
||||
tree c, *pc, mapper_name = NULL_TREE;
|
||||
|
||||
for (pc = &clauses, c = clauses; c; c = *pc)
|
||||
{
|
||||
bool using_mapper = false;
|
||||
|
||||
switch (OMP_CLAUSE_CODE (c))
|
||||
{
|
||||
case OMP_CLAUSE_MAP:
|
||||
{
|
||||
tree t = OMP_CLAUSE_DECL (c);
|
||||
tree type = NULL_TREE;
|
||||
bool nonunit_array_with_mapper = false;
|
||||
|
||||
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME
|
||||
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POP_MAPPER_NAME)
|
||||
{
|
||||
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME)
|
||||
mapper_name = OMP_CLAUSE_DECL (c);
|
||||
else
|
||||
mapper_name = NULL_TREE;
|
||||
pc = &OMP_CLAUSE_CHAIN (c);
|
||||
continue;
|
||||
}
|
||||
|
||||
if (TREE_CODE (t) == OMP_ARRAY_SECTION)
|
||||
{
|
||||
location_t loc = OMP_CLAUSE_LOCATION (c);
|
||||
tree t2 = lang_hooks.decls.omp_map_array_section (loc, t);
|
||||
|
||||
if (t2 == t)
|
||||
{
|
||||
/* !!! Array sections of size >1 with mappers for elements
|
||||
are hard to support. Do something here. */
|
||||
nonunit_array_with_mapper = true;
|
||||
type = TREE_TYPE (TREE_TYPE (t));
|
||||
}
|
||||
else
|
||||
{
|
||||
t = t2;
|
||||
type = TREE_TYPE (t);
|
||||
}
|
||||
}
|
||||
else
|
||||
type = TREE_TYPE (t);
|
||||
|
||||
if (type == NULL_TREE || type == error_mark_node)
|
||||
{
|
||||
pc = &OMP_CLAUSE_CHAIN (c);
|
||||
continue;
|
||||
}
|
||||
|
||||
enum gomp_map_kind kind = OMP_CLAUSE_MAP_KIND (c);
|
||||
if (kind == GOMP_MAP_UNSET)
|
||||
kind = GOMP_MAP_TOFROM;
|
||||
|
||||
type = TYPE_MAIN_VARIANT (type);
|
||||
|
||||
tree mapper_fn
|
||||
= lang_hooks.decls.omp_mapper_lookup (mapper_name, type);
|
||||
|
||||
if (mapper_fn && nonunit_array_with_mapper)
|
||||
{
|
||||
sorry ("user-defined mapper with non-unit length "
|
||||
"array section");
|
||||
using_mapper = true;
|
||||
}
|
||||
else if (mapper_fn)
|
||||
{
|
||||
tree mapper
|
||||
= lang_hooks.decls.omp_extract_mapper_directive (mapper_fn);
|
||||
pc = omp_instantiate_mapper (pc, mapper, t, kind);
|
||||
using_mapper = true;
|
||||
}
|
||||
else if (mapper_name)
|
||||
{
|
||||
error ("mapper %qE not found for type %qT", mapper_name, type);
|
||||
using_mapper = true;
|
||||
}
|
||||
}
|
||||
break;
|
||||
|
||||
default:
|
||||
;
|
||||
}
|
||||
|
||||
if (using_mapper)
|
||||
*pc = OMP_CLAUSE_CHAIN (c);
|
||||
else
|
||||
pc = &OMP_CLAUSE_CHAIN (c);
|
||||
}
|
||||
|
||||
return clauses;
|
||||
}
|
||||
|
||||
const struct c_omp_directive c_omp_directives[] = {
|
||||
/* Keep this alphabetically sorted by the first word. Non-null second/third
|
||||
if any should precede null ones. */
|
||||
|
|
|
|||
|
|
@ -3529,6 +3529,9 @@ reduced_constant_expression_p (tree t)
|
|||
/* Even if we can't lower this yet, it's constant. */
|
||||
return true;
|
||||
|
||||
case OMP_DECLARE_MAPPER:
|
||||
return true;
|
||||
|
||||
case CONSTRUCTOR:
|
||||
/* And we need to handle PTRMEM_CST wrapped in a CONSTRUCTOR. */
|
||||
tree field;
|
||||
|
|
@ -7838,6 +7841,7 @@ cxx_eval_constant_expression (const constexpr_ctx *ctx, tree t,
|
|||
case LABEL_EXPR:
|
||||
case CASE_LABEL_EXPR:
|
||||
case PREDICT_EXPR:
|
||||
case OMP_DECLARE_MAPPER:
|
||||
return t;
|
||||
|
||||
case PARM_DECL:
|
||||
|
|
@ -10532,6 +10536,11 @@ potential_constant_expression_1 (tree t, bool want_rval, bool strict, bool now,
|
|||
"expression", t);
|
||||
return false;
|
||||
|
||||
case OMP_DECLARE_MAPPER:
|
||||
/* This can be used to initialize VAR_DECLs: it's treated as a magic
|
||||
constant. */
|
||||
return true;
|
||||
|
||||
case ASM_EXPR:
|
||||
if (flags & tf_error)
|
||||
inline_asm_in_constexpr_error (loc, fundef_p);
|
||||
|
|
|
|||
|
|
@ -2809,6 +2809,12 @@ cxx_omp_finish_clause (tree c, gimple_seq *, bool /* openacc */)
|
|||
}
|
||||
}
|
||||
|
||||
tree
|
||||
cxx_omp_finish_mapper_clauses (tree clauses)
|
||||
{
|
||||
return finish_omp_clauses (clauses, C_ORT_OMP);
|
||||
}
|
||||
|
||||
/* Return true if DECL's DECL_VALUE_EXPR (if any) should be
|
||||
disregarded in OpenMP construct, because it is going to be
|
||||
remapped during OpenMP lowering. SHARED is true if DECL
|
||||
|
|
|
|||
|
|
@ -190,6 +190,15 @@ static const scoped_attribute_specs *const cp_objcp_attribute_table[] =
|
|||
#define LANG_HOOKS_OMP_CLAUSE_DTOR cxx_omp_clause_dtor
|
||||
#undef LANG_HOOKS_OMP_FINISH_CLAUSE
|
||||
#define LANG_HOOKS_OMP_FINISH_CLAUSE cxx_omp_finish_clause
|
||||
#undef LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES
|
||||
#define LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES cxx_omp_finish_mapper_clauses
|
||||
#undef LANG_HOOKS_OMP_MAPPER_LOOKUP
|
||||
#define LANG_HOOKS_OMP_MAPPER_LOOKUP cxx_omp_mapper_lookup
|
||||
#undef LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE
|
||||
#define LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE \
|
||||
cxx_omp_extract_mapper_directive
|
||||
#undef LANG_HOOKS_OMP_MAP_ARRAY_SECTION
|
||||
#define LANG_HOOKS_OMP_MAP_ARRAY_SECTION cxx_omp_map_array_section
|
||||
#undef LANG_HOOKS_OMP_PRIVATIZE_BY_REFERENCE
|
||||
#define LANG_HOOKS_OMP_PRIVATIZE_BY_REFERENCE cxx_omp_privatize_by_reference
|
||||
#undef LANG_HOOKS_OMP_DISREGARD_VALUE_EXPR
|
||||
|
|
|
|||
|
|
@ -2975,7 +2975,10 @@ struct GTY(()) lang_decl_base {
|
|||
|
||||
unsigned module_keyed_decls_p : 1; /* has keys, applies to all decls */
|
||||
|
||||
/* 11 spare bits. */
|
||||
/* VAR_DECL being used to represent an OpenMP declared mapper. */
|
||||
unsigned omp_declare_mapper_p : 1;
|
||||
|
||||
/* 10 spare bits. */
|
||||
};
|
||||
|
||||
/* True for DECL codes which have template info and access. */
|
||||
|
|
@ -4530,6 +4533,11 @@ get_vec_init_expr (tree t)
|
|||
#define DECL_OMP_DECLARE_REDUCTION_P(NODE) \
|
||||
(LANG_DECL_FN_CHECK (DECL_COMMON_CHECK (NODE))->omp_declare_reduction_p)
|
||||
|
||||
/* Nonzero if NODE is an artificial FUNCTION_DECL for
|
||||
#pragma omp declare mapper. */
|
||||
#define DECL_OMP_DECLARE_MAPPER_P(NODE) \
|
||||
(DECL_LANG_SPECIFIC (VAR_DECL_CHECK (NODE))->u.base.omp_declare_mapper_p)
|
||||
|
||||
/* Nonzero if DECL has been declared threadprivate by
|
||||
#pragma omp threadprivate. */
|
||||
#define CP_DECL_THREADPRIVATE_P(DECL) \
|
||||
|
|
@ -8096,11 +8104,14 @@ extern tree finish_qualified_id_expr (tree, tree, bool, bool,
|
|||
extern void simplify_aggr_init_expr (tree *);
|
||||
extern void finalize_nrv (tree, tree);
|
||||
extern tree omp_reduction_id (enum tree_code, tree, tree);
|
||||
extern tree omp_mapper_id (tree, tree);
|
||||
extern tree cp_remove_omp_priv_cleanup_stmt (tree *, int *, void *);
|
||||
extern bool cp_check_omp_declare_reduction (tree);
|
||||
extern bool cp_check_omp_declare_mapper (tree);
|
||||
extern void finish_omp_declare_simd_methods (tree);
|
||||
extern tree cp_finish_omp_init_prefer_type (tree);
|
||||
extern tree finish_omp_clauses (tree, enum c_omp_region_type);
|
||||
extern tree omp_instantiate_mappers (tree);
|
||||
extern tree push_omp_privatization_clauses (bool);
|
||||
extern void pop_omp_privatization_clauses (tree);
|
||||
extern void save_omp_privatization_clauses (vec<tree> &);
|
||||
|
|
@ -8696,6 +8707,10 @@ extern tree cxx_omp_clause_copy_ctor (tree, tree, tree);
|
|||
extern tree cxx_omp_clause_assign_op (tree, tree, tree);
|
||||
extern tree cxx_omp_clause_dtor (tree, tree);
|
||||
extern void cxx_omp_finish_clause (tree, gimple_seq *, bool);
|
||||
extern tree cxx_omp_finish_mapper_clauses (tree);
|
||||
extern tree cxx_omp_mapper_lookup (tree, tree);
|
||||
extern tree cxx_omp_extract_mapper_directive (tree);
|
||||
extern tree cxx_omp_map_array_section (location_t, tree);
|
||||
extern bool cxx_omp_privatize_by_reference (const_tree);
|
||||
extern bool cxx_omp_disregard_value_expr (tree, bool);
|
||||
extern void cp_fold_function (tree);
|
||||
|
|
|
|||
|
|
@ -7877,6 +7877,12 @@ check_initializer (tree decl, tree init, int flags, vec<tree, va_gc> **cleanups)
|
|||
}
|
||||
else if (!init && DECL_REALLY_EXTERN (decl))
|
||||
;
|
||||
else if (flag_openmp
|
||||
&& VAR_P (decl)
|
||||
&& DECL_LANG_SPECIFIC (decl)
|
||||
&& DECL_OMP_DECLARE_MAPPER_P (decl)
|
||||
&& TREE_CODE (init) == OMP_DECLARE_MAPPER)
|
||||
return NULL_TREE;
|
||||
else if (init || type_build_ctor_call (type)
|
||||
|| TYPE_REF_P (type))
|
||||
{
|
||||
|
|
@ -9197,14 +9203,23 @@ cp_finish_decl (tree decl, tree init, bool init_const_expr_p,
|
|||
varpool_node::get_create (decl);
|
||||
}
|
||||
|
||||
if (flag_openmp
|
||||
&& VAR_P (decl)
|
||||
&& DECL_LANG_SPECIFIC (decl)
|
||||
&& DECL_OMP_DECLARE_MAPPER_P (decl)
|
||||
&& init)
|
||||
{
|
||||
gcc_assert (TREE_CODE (init) == OMP_DECLARE_MAPPER);
|
||||
DECL_INITIAL (decl) = init;
|
||||
}
|
||||
/* Convert the initializer to the type of DECL, if we have not
|
||||
already initialized DECL. */
|
||||
if (!DECL_INITIALIZED_P (decl)
|
||||
/* If !DECL_EXTERNAL then DECL is being defined. In the
|
||||
case of a static data member initialized inside the
|
||||
class-specifier, there can be an initializer even if DECL
|
||||
is *not* defined. */
|
||||
&& (!DECL_EXTERNAL (decl) || init))
|
||||
else if (!DECL_INITIALIZED_P (decl)
|
||||
/* If !DECL_EXTERNAL then DECL is being defined. In the
|
||||
case of a static data member initialized inside the
|
||||
class-specifier, there can be an initializer even if DECL
|
||||
is *not* defined. */
|
||||
&& (!DECL_EXTERNAL (decl) || init))
|
||||
{
|
||||
cleanups = make_tree_vector ();
|
||||
init = check_initializer (decl, init, flags, &cleanups);
|
||||
|
|
|
|||
|
|
@ -6408,12 +6408,17 @@ mark_used (tree decl, tsubst_flags_t complain /* = tf_warning_or_error */)
|
|||
|
||||
/* If DECL has a deduced return type, we need to instantiate it now to
|
||||
find out its type. For OpenMP user defined reductions, we need them
|
||||
instantiated for reduction clauses which inline them by hand directly. */
|
||||
instantiated for reduction clauses which inline them by hand directly.
|
||||
OpenMP declared mappers are used implicitly so must be instantiated
|
||||
before they can be detected. */
|
||||
if (undeduced_auto_decl (decl)
|
||||
|| (VAR_P (decl)
|
||||
&& VAR_HAD_UNKNOWN_BOUND (decl))
|
||||
|| (TREE_CODE (decl) == FUNCTION_DECL
|
||||
&& DECL_OMP_DECLARE_REDUCTION_P (decl)))
|
||||
&& DECL_OMP_DECLARE_REDUCTION_P (decl))
|
||||
|| (TREE_CODE (decl) == VAR_DECL
|
||||
&& DECL_LANG_SPECIFIC (decl)
|
||||
&& DECL_OMP_DECLARE_MAPPER_P (decl)))
|
||||
maybe_instantiate_decl (decl);
|
||||
|
||||
if (!decl_dependent_p (decl)
|
||||
|
|
|
|||
|
|
@ -1282,12 +1282,37 @@ dump_global_iord (cxx_pretty_printer *pp, tree t)
|
|||
pp_printf (pp, p, DECL_SOURCE_FILE (t));
|
||||
}
|
||||
|
||||
/* Write a representation of OpenMP "declare mapper" T to PP in a manner
|
||||
suitable for error messages. */
|
||||
|
||||
static void
|
||||
dump_omp_declare_mapper (cxx_pretty_printer *pp, tree t, int flags)
|
||||
{
|
||||
pp_string (pp, "#pragma omp declare mapper");
|
||||
if (t == NULL_TREE || t == error_mark_node)
|
||||
return;
|
||||
pp_space (pp);
|
||||
pp_cxx_left_paren (pp);
|
||||
if (OMP_DECLARE_MAPPER_ID (t))
|
||||
{
|
||||
pp_cxx_tree_identifier (pp, OMP_DECLARE_MAPPER_ID (t));
|
||||
pp_colon (pp);
|
||||
}
|
||||
dump_type (pp, TREE_TYPE (t), flags);
|
||||
pp_cxx_right_paren (pp);
|
||||
}
|
||||
|
||||
static void
|
||||
dump_simple_decl (cxx_pretty_printer *pp, tree t, tree type, int flags)
|
||||
{
|
||||
if (VAR_P (t) && DECL_NTTP_OBJECT_P (t))
|
||||
return dump_expr (pp, DECL_INITIAL (t), flags);
|
||||
|
||||
if (TREE_CODE (t) == VAR_DECL
|
||||
&& DECL_LANG_SPECIFIC (t)
|
||||
&& DECL_OMP_DECLARE_MAPPER_P (t))
|
||||
return dump_omp_declare_mapper (pp, DECL_INITIAL (t), flags);
|
||||
|
||||
if (flags & TFF_DECL_SPECIFIERS)
|
||||
{
|
||||
if (concept_definition_p (t))
|
||||
|
|
|
|||
299
gcc/cp/parser.cc
299
gcc/cp/parser.cc
|
|
@ -42424,13 +42424,13 @@ cp_parser_omp_clause_from_to (cp_parser *parser, enum omp_clause_code kind,
|
|||
map ( [map-type-modifier[,] ...] map-kind: variable-list )
|
||||
|
||||
map-type-modifier:
|
||||
always | close */
|
||||
always | close | mapper ( mapper-name ) */
|
||||
|
||||
static tree
|
||||
cp_parser_omp_clause_map (cp_parser *parser, tree list)
|
||||
cp_parser_omp_clause_map (cp_parser *parser, tree list, bool declare_mapper_p)
|
||||
{
|
||||
tree nlist, c;
|
||||
enum gomp_map_kind kind = GOMP_MAP_TOFROM;
|
||||
enum gomp_map_kind kind = declare_mapper_p ? GOMP_MAP_UNSET : GOMP_MAP_TOFROM;
|
||||
|
||||
if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
|
||||
return list;
|
||||
|
|
@ -42448,12 +42448,17 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
|
|||
|
||||
if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type == CPP_COMMA)
|
||||
pos++;
|
||||
else if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type
|
||||
== CPP_OPEN_PAREN)
|
||||
pos = cp_parser_skip_balanced_tokens (parser, pos + 1);
|
||||
pos++;
|
||||
}
|
||||
|
||||
bool always_modifier = false;
|
||||
bool close_modifier = false;
|
||||
bool present_modifier = false;
|
||||
bool mapper_modifier = false;
|
||||
tree mapper_name = NULL_TREE;
|
||||
for (int pos = 1; pos < map_kind_pos; ++pos)
|
||||
{
|
||||
cp_token *tok = cp_lexer_peek_token (parser->lexer);
|
||||
|
|
@ -42476,6 +42481,7 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
|
|||
return list;
|
||||
}
|
||||
always_modifier = true;
|
||||
cp_lexer_consume_token (parser->lexer);
|
||||
}
|
||||
else if (strcmp ("close", p) == 0)
|
||||
{
|
||||
|
|
@ -42489,6 +42495,77 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
|
|||
return list;
|
||||
}
|
||||
close_modifier = true;
|
||||
cp_lexer_consume_token (parser->lexer);
|
||||
}
|
||||
else if (strcmp ("mapper", p) == 0)
|
||||
{
|
||||
cp_lexer_consume_token (parser->lexer);
|
||||
|
||||
matching_parens parens;
|
||||
if (parens.require_open (parser))
|
||||
{
|
||||
if (mapper_modifier)
|
||||
{
|
||||
cp_parser_error (parser, "too many %<mapper%> modifiers");
|
||||
/* Assume it's a well-formed mapper modifier, even if it
|
||||
seems to be in the wrong place. */
|
||||
cp_lexer_consume_token (parser->lexer);
|
||||
parens.require_close (parser);
|
||||
cp_parser_skip_to_closing_parenthesis (parser,
|
||||
/*recovering=*/true,
|
||||
/*or_comma=*/false,
|
||||
/*consume_paren=*/
|
||||
true);
|
||||
return list;
|
||||
}
|
||||
|
||||
tok = cp_lexer_peek_token (parser->lexer);
|
||||
switch (tok->type)
|
||||
{
|
||||
case CPP_NAME:
|
||||
{
|
||||
cp_expr e = cp_parser_identifier (parser);
|
||||
if (e != error_mark_node)
|
||||
mapper_name = e;
|
||||
else
|
||||
goto err;
|
||||
if (declare_mapper_p)
|
||||
{
|
||||
error_at (e.get_location (),
|
||||
"in %<declare mapper%> directives, parameter "
|
||||
"to %<mapper%> modifier must be %<default%>");
|
||||
}
|
||||
}
|
||||
break;
|
||||
|
||||
case CPP_KEYWORD:
|
||||
if (tok->keyword == RID_DEFAULT)
|
||||
{
|
||||
cp_lexer_consume_token (parser->lexer);
|
||||
break;
|
||||
}
|
||||
/* Fallthrough. */
|
||||
|
||||
default:
|
||||
err:
|
||||
cp_parser_error (parser,
|
||||
"expected identifier or %<default%>");
|
||||
return list;
|
||||
}
|
||||
|
||||
if (!parens.require_close (parser))
|
||||
{
|
||||
cp_parser_skip_to_closing_parenthesis (parser,
|
||||
/*recovering=*/true,
|
||||
/*or_comma=*/false,
|
||||
/*consume_paren=*/
|
||||
true);
|
||||
return list;
|
||||
}
|
||||
|
||||
mapper_modifier = true;
|
||||
pos += 3;
|
||||
}
|
||||
}
|
||||
else if (strcmp ("present", p) == 0)
|
||||
{
|
||||
|
|
@ -42502,19 +42579,19 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
|
|||
return list;
|
||||
}
|
||||
present_modifier = true;
|
||||
}
|
||||
cp_lexer_consume_token (parser->lexer);
|
||||
}
|
||||
else
|
||||
{
|
||||
cp_parser_error (parser, "%<map%> clause with map-type modifier other"
|
||||
" than %<always%>, %<close%> or %<present%>");
|
||||
cp_parser_error (parser, "%<map%> clause with map-type modifier "
|
||||
"other than %<always%>, %<close%>, "
|
||||
"%<mapper%> or %<present%>");
|
||||
cp_parser_skip_to_closing_parenthesis (parser,
|
||||
/*recovering=*/true,
|
||||
/*or_comma=*/false,
|
||||
/*consume_paren=*/true);
|
||||
return list;
|
||||
}
|
||||
|
||||
cp_lexer_consume_token (parser->lexer);
|
||||
}
|
||||
|
||||
if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)
|
||||
|
|
@ -42570,8 +42647,30 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
|
|||
NULL, true);
|
||||
finish_scope ();
|
||||
|
||||
tree last_new = NULL_TREE;
|
||||
|
||||
for (c = nlist; c != list; c = OMP_CLAUSE_CHAIN (c))
|
||||
OMP_CLAUSE_SET_MAP_KIND (c, kind);
|
||||
{
|
||||
OMP_CLAUSE_SET_MAP_KIND (c, kind);
|
||||
last_new = c;
|
||||
}
|
||||
|
||||
if (mapper_name)
|
||||
{
|
||||
tree name = build_omp_clause (input_location, OMP_CLAUSE_MAP);
|
||||
OMP_CLAUSE_SET_MAP_KIND (name, GOMP_MAP_PUSH_MAPPER_NAME);
|
||||
OMP_CLAUSE_DECL (name) = mapper_name;
|
||||
OMP_CLAUSE_CHAIN (name) = nlist;
|
||||
nlist = name;
|
||||
|
||||
gcc_assert (last_new);
|
||||
|
||||
name = build_omp_clause (input_location, OMP_CLAUSE_MAP);
|
||||
OMP_CLAUSE_SET_MAP_KIND (name, GOMP_MAP_POP_MAPPER_NAME);
|
||||
OMP_CLAUSE_DECL (name) = null_pointer_node;
|
||||
OMP_CLAUSE_CHAIN (name) = OMP_CLAUSE_CHAIN (last_new);
|
||||
OMP_CLAUSE_CHAIN (last_new) = name;
|
||||
}
|
||||
|
||||
return nlist;
|
||||
}
|
||||
|
|
@ -43899,7 +43998,8 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask,
|
|||
c_name = "detach";
|
||||
break;
|
||||
case PRAGMA_OMP_CLAUSE_MAP:
|
||||
clauses = cp_parser_omp_clause_map (parser, clauses);
|
||||
clauses = cp_parser_omp_clause_map (parser, clauses,
|
||||
/*mapper=*/false);
|
||||
c_name = "map";
|
||||
break;
|
||||
case PRAGMA_OMP_CLAUSE_DEVICE:
|
||||
|
|
@ -48905,6 +49005,8 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
|
|||
OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c);
|
||||
OMP_CLAUSE_CHAIN (c) = nc;
|
||||
}
|
||||
if (!processing_template_decl)
|
||||
clauses = c_omp_instantiate_mappers (clauses);
|
||||
clauses = finish_omp_clauses (clauses, C_ORT_OMP_TARGET);
|
||||
|
||||
c_omp_adjust_map_clauses (clauses, true);
|
||||
|
|
@ -52179,6 +52281,175 @@ cp_parser_omp_declare_reduction (cp_parser *parser, cp_token *pragma_tok,
|
|||
obstack_free (&declarator_obstack, p);
|
||||
}
|
||||
|
||||
/* OpenMP 5.0
|
||||
#pragma omp declare mapper([mapper-identifier:]type var) \
|
||||
[clause[[,] clause] ... ] new-line */
|
||||
|
||||
static void
|
||||
cp_parser_omp_declare_mapper (cp_parser *parser, cp_token *pragma_tok,
|
||||
enum pragma_context)
|
||||
{
|
||||
cp_token *token = NULL;
|
||||
tree type = NULL_TREE, vardecl = NULL_TREE, block = NULL_TREE;
|
||||
bool block_scope = false;
|
||||
/* Don't create location wrapper nodes within "declare mapper"
|
||||
directives. */
|
||||
auto_suppress_location_wrappers sentinel;
|
||||
tree mapper_name = NULL_TREE;
|
||||
tree mapper_id, id, placeholder, mapper, maplist = NULL_TREE;
|
||||
|
||||
if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
|
||||
goto fail;
|
||||
|
||||
if (current_function_decl)
|
||||
block_scope = true;
|
||||
|
||||
token = cp_lexer_peek_token (parser->lexer);
|
||||
|
||||
if (cp_lexer_nth_token_is (parser->lexer, 2, CPP_COLON))
|
||||
{
|
||||
switch (token->type)
|
||||
{
|
||||
case CPP_NAME:
|
||||
{
|
||||
cp_expr e = cp_parser_identifier (parser);
|
||||
if (e != error_mark_node)
|
||||
mapper_name = e;
|
||||
else
|
||||
goto fail;
|
||||
}
|
||||
break;
|
||||
|
||||
case CPP_KEYWORD:
|
||||
if (token->keyword == RID_DEFAULT)
|
||||
{
|
||||
mapper_name = NULL_TREE;
|
||||
cp_lexer_consume_token (parser->lexer);
|
||||
break;
|
||||
}
|
||||
/* Fallthrough. */
|
||||
|
||||
default:
|
||||
cp_parser_error (parser, "expected identifier or %<default%>");
|
||||
}
|
||||
|
||||
if (!cp_parser_require (parser, CPP_COLON, RT_COLON))
|
||||
goto fail;
|
||||
}
|
||||
|
||||
{
|
||||
const char *saved_message = parser->type_definition_forbidden_message;
|
||||
parser->type_definition_forbidden_message
|
||||
= G_("types may not be defined within %<declare mapper%>");
|
||||
type_id_in_expr_sentinel s (parser);
|
||||
type = cp_parser_type_id (parser);
|
||||
parser->type_definition_forbidden_message = saved_message;
|
||||
}
|
||||
|
||||
if (type == error_mark_node)
|
||||
goto fail;
|
||||
if (dependent_type_p (type))
|
||||
mapper_id = omp_mapper_id (mapper_name, NULL_TREE);
|
||||
else
|
||||
mapper_id = omp_mapper_id (mapper_name, type);
|
||||
|
||||
vardecl = build_lang_decl (VAR_DECL, mapper_id, type);
|
||||
DECL_ARTIFICIAL (vardecl) = 1;
|
||||
TREE_STATIC (vardecl) = 1;
|
||||
TREE_PUBLIC (vardecl) = 0;
|
||||
DECL_EXTERNAL (vardecl) = 0;
|
||||
DECL_DECLARED_CONSTEXPR_P (vardecl) = 1;
|
||||
DECL_INITIALIZED_BY_CONSTANT_EXPRESSION_P (vardecl) = 1;
|
||||
DECL_OMP_DECLARE_MAPPER_P (vardecl) = 1;
|
||||
|
||||
keep_next_level (true);
|
||||
block = begin_omp_structured_block ();
|
||||
|
||||
if (block_scope)
|
||||
DECL_CONTEXT (vardecl) = current_function_decl;
|
||||
else if (current_class_type)
|
||||
DECL_CONTEXT (vardecl) = current_class_type;
|
||||
else
|
||||
DECL_CONTEXT (vardecl) = current_namespace;
|
||||
|
||||
if (processing_template_decl)
|
||||
vardecl = push_template_decl (vardecl);
|
||||
|
||||
if ((id = cp_parser_declarator_id (parser, false)) == error_mark_node)
|
||||
goto fail;
|
||||
|
||||
if (!cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN))
|
||||
{
|
||||
finish_omp_structured_block (block);
|
||||
goto fail;
|
||||
}
|
||||
|
||||
placeholder = build_lang_decl (VAR_DECL, id, type);
|
||||
DECL_CONTEXT (placeholder) = DECL_CONTEXT (vardecl);
|
||||
if (processing_template_decl)
|
||||
placeholder = push_template_decl (placeholder);
|
||||
pushdecl (placeholder);
|
||||
cp_finish_decl (placeholder, NULL_TREE, 0, NULL_TREE, 0);
|
||||
DECL_ARTIFICIAL (placeholder) = 1;
|
||||
TREE_USED (placeholder) = 1;
|
||||
|
||||
while (cp_lexer_next_token_is_not (parser->lexer, CPP_PRAGMA_EOL))
|
||||
{
|
||||
pragma_omp_clause c_kind = cp_parser_omp_clause_name (parser);
|
||||
if (c_kind != PRAGMA_OMP_CLAUSE_MAP)
|
||||
{
|
||||
if (c_kind != PRAGMA_OMP_CLAUSE_NONE)
|
||||
cp_parser_error (parser, "unexpected clause");
|
||||
finish_omp_structured_block (block);
|
||||
goto fail;
|
||||
}
|
||||
maplist = cp_parser_omp_clause_map (parser, maplist, /*mapper=*/true);
|
||||
if (maplist == NULL_TREE)
|
||||
break;
|
||||
}
|
||||
|
||||
if (maplist == NULL_TREE)
|
||||
{
|
||||
cp_parser_error (parser, "missing %<map%> clause");
|
||||
finish_omp_structured_block (block);
|
||||
goto fail;
|
||||
}
|
||||
|
||||
mapper = make_node (OMP_DECLARE_MAPPER);
|
||||
TREE_TYPE (mapper) = type;
|
||||
OMP_DECLARE_MAPPER_ID (mapper) = mapper_name;
|
||||
OMP_DECLARE_MAPPER_DECL (mapper) = placeholder;
|
||||
OMP_DECLARE_MAPPER_CLAUSES (mapper) = maplist;
|
||||
|
||||
finish_omp_structured_block (block);
|
||||
|
||||
DECL_INITIAL (vardecl) = mapper;
|
||||
|
||||
if (current_class_type)
|
||||
{
|
||||
if (processing_template_decl)
|
||||
{
|
||||
retrofit_lang_decl (vardecl);
|
||||
SET_DECL_VAR_DECLARED_INLINE_P (vardecl);
|
||||
}
|
||||
finish_static_data_member_decl (vardecl, mapper,
|
||||
/*init_const_expr_p=*/true, NULL_TREE, 0);
|
||||
finish_member_declaration (vardecl);
|
||||
}
|
||||
else if (processing_template_decl && block_scope)
|
||||
add_decl_expr (vardecl);
|
||||
else
|
||||
pushdecl (vardecl);
|
||||
|
||||
cp_check_omp_declare_mapper (vardecl);
|
||||
|
||||
cp_parser_require_pragma_eol (parser, pragma_tok);
|
||||
return;
|
||||
|
||||
fail:
|
||||
cp_parser_skip_to_pragma_eol (parser, pragma_tok);
|
||||
}
|
||||
|
||||
/* OpenMP 4.0
|
||||
#pragma omp declare simd declare-simd-clauses[optseq] new-line
|
||||
#pragma omp declare reduction (reduction-id : typename-list : expression) \
|
||||
|
|
@ -52223,6 +52494,12 @@ cp_parser_omp_declare (cp_parser *parser, cp_token *pragma_tok,
|
|||
context);
|
||||
return false;
|
||||
}
|
||||
if (strcmp (p, "mapper") == 0)
|
||||
{
|
||||
cp_lexer_consume_token (parser->lexer);
|
||||
cp_parser_omp_declare_mapper (parser, pragma_tok, context);
|
||||
return false;
|
||||
}
|
||||
if (!flag_openmp) /* flag_openmp_simd */
|
||||
{
|
||||
cp_parser_skip_to_pragma_eol (parser, pragma_tok);
|
||||
|
|
@ -52236,7 +52513,7 @@ cp_parser_omp_declare (cp_parser *parser, cp_token *pragma_tok,
|
|||
}
|
||||
}
|
||||
cp_parser_error (parser, "expected %<simd%>, %<reduction%>, "
|
||||
"%<target%> or %<variant%>");
|
||||
"%<target%>, %<mapper%> or %<variant%>");
|
||||
cp_parser_require_pragma_eol (parser, pragma_tok);
|
||||
return false;
|
||||
}
|
||||
|
|
|
|||
53
gcc/cp/pt.cc
53
gcc/cp/pt.cc
|
|
@ -15989,7 +15989,10 @@ tsubst_decl (tree t, tree args, tsubst_flags_t complain,
|
|||
= remove_attribute ("visibility", DECL_ATTRIBUTES (r));
|
||||
}
|
||||
determine_visibility (r);
|
||||
if ((!local_p || TREE_STATIC (t)) && DECL_SECTION_NAME (t))
|
||||
if ((!local_p || TREE_STATIC (t))
|
||||
&& !(flag_openmp && DECL_LANG_SPECIFIC (t)
|
||||
&& DECL_OMP_DECLARE_MAPPER_P (t))
|
||||
&& DECL_SECTION_NAME (t))
|
||||
set_decl_section_name (r, t);
|
||||
}
|
||||
|
||||
|
|
@ -16041,6 +16044,13 @@ tsubst_decl (tree t, tree args, tsubst_flags_t complain,
|
|||
SET_TYPE_STRUCTURAL_EQUALITY (TREE_TYPE (r));
|
||||
}
|
||||
|
||||
if (flag_openmp
|
||||
&& VAR_P (t)
|
||||
&& DECL_LANG_SPECIFIC (t)
|
||||
&& DECL_OMP_DECLARE_MAPPER_P (t)
|
||||
&& strchr (IDENTIFIER_POINTER (DECL_NAME (t)), '~') == NULL)
|
||||
DECL_NAME (r) = omp_mapper_id (DECL_NAME (t), TREE_TYPE (r));
|
||||
|
||||
layout_decl (r, 0);
|
||||
}
|
||||
break;
|
||||
|
|
@ -18313,8 +18323,10 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort,
|
|||
}
|
||||
|
||||
new_clauses = nreverse (new_clauses);
|
||||
if (ort != C_ORT_OMP_DECLARE_SIMD)
|
||||
if (ort != C_ORT_OMP_DECLARE_SIMD && ort != C_ORT_OMP_DECLARE_MAPPER)
|
||||
{
|
||||
if (ort == C_ORT_OMP_TARGET)
|
||||
new_clauses = c_omp_instantiate_mappers (new_clauses);
|
||||
new_clauses = finish_omp_clauses (new_clauses, ort);
|
||||
if (linear_no_step)
|
||||
for (nc = new_clauses; nc; nc = OMP_CLAUSE_CHAIN (nc))
|
||||
|
|
@ -20034,6 +20046,22 @@ tsubst_stmt (tree t, tree args, tsubst_flags_t complain, tree in_decl)
|
|||
break;
|
||||
}
|
||||
|
||||
case OMP_DECLARE_MAPPER:
|
||||
{
|
||||
t = copy_node (t);
|
||||
|
||||
tree decl = OMP_DECLARE_MAPPER_DECL (t);
|
||||
decl = tsubst (decl, args, complain, in_decl);
|
||||
tree type = tsubst (TREE_TYPE (t), args, complain, in_decl);
|
||||
tree clauses = OMP_DECLARE_MAPPER_CLAUSES (t);
|
||||
clauses = tsubst_omp_clauses (clauses, C_ORT_OMP_DECLARE_MAPPER, args,
|
||||
complain, in_decl);
|
||||
TREE_TYPE (t) = type;
|
||||
OMP_DECLARE_MAPPER_DECL (t) = decl;
|
||||
OMP_DECLARE_MAPPER_CLAUSES (t) = clauses;
|
||||
RETURN (t);
|
||||
}
|
||||
|
||||
case TRANSACTION_EXPR:
|
||||
{
|
||||
int flags = 0;
|
||||
|
|
@ -21093,6 +21121,23 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl)
|
|||
RETURN (build_omp_array_section (EXPR_LOCATION (t), op0, op1, op2));
|
||||
}
|
||||
|
||||
case OMP_DECLARE_MAPPER:
|
||||
{
|
||||
t = copy_node (t);
|
||||
|
||||
tree decl = OMP_DECLARE_MAPPER_DECL (t);
|
||||
DECL_OMP_DECLARE_MAPPER_P (decl) = 1;
|
||||
decl = tsubst (decl, args, complain, in_decl);
|
||||
tree type = tsubst (TREE_TYPE (t), args, complain, in_decl);
|
||||
tree clauses = OMP_DECLARE_MAPPER_CLAUSES (t);
|
||||
clauses = tsubst_omp_clauses (clauses, C_ORT_OMP_DECLARE_MAPPER, args,
|
||||
complain, in_decl);
|
||||
TREE_TYPE (t) = type;
|
||||
OMP_DECLARE_MAPPER_DECL (t) = decl;
|
||||
OMP_DECLARE_MAPPER_CLAUSES (t) = clauses;
|
||||
RETURN (t);
|
||||
}
|
||||
|
||||
case SIZEOF_EXPR:
|
||||
if (PACK_EXPANSION_P (TREE_OPERAND (t, 0))
|
||||
|| ARGUMENT_PACK_P (TREE_OPERAND (t, 0)))
|
||||
|
|
@ -28036,7 +28081,9 @@ instantiate_decl (tree d, bool defer_ok, bool expl_inst_class_mem_p)
|
|||
|| (external_p && VAR_P (d))
|
||||
/* Handle here a deleted function too, avoid generating
|
||||
its body (c++/61080). */
|
||||
|| deleted_p)
|
||||
|| deleted_p
|
||||
/* We need the initializer for an OpenMP declare mapper. */
|
||||
|| (VAR_P (d) && DECL_LANG_SPECIFIC (d) && DECL_OMP_DECLARE_MAPPER_P (d)))
|
||||
{
|
||||
/* The definition of the static data member is now required so
|
||||
we must substitute the initializer. */
|
||||
|
|
|
|||
|
|
@ -45,6 +45,7 @@ along with GCC; see the file COPYING3. If not see
|
|||
#include "gomp-constants.h"
|
||||
#include "predict.h"
|
||||
#include "memmodel.h"
|
||||
#include "gimplify.h"
|
||||
|
||||
/* There routines provide a modular interface to perform many parsing
|
||||
operations. They may therefore be used during actual parsing, or
|
||||
|
|
@ -6740,6 +6741,97 @@ omp_reduction_lookup (location_t loc, tree id, tree type, tree *baselinkp,
|
|||
return id;
|
||||
}
|
||||
|
||||
/* Return identifier to look up for omp declare mapper. */
|
||||
|
||||
tree
|
||||
omp_mapper_id (tree mapper_id, tree type)
|
||||
{
|
||||
const char *p = NULL;
|
||||
const char *m = NULL;
|
||||
|
||||
if (mapper_id == NULL_TREE)
|
||||
p = "";
|
||||
else if (TREE_CODE (mapper_id) == IDENTIFIER_NODE)
|
||||
p = IDENTIFIER_POINTER (mapper_id);
|
||||
else
|
||||
return error_mark_node;
|
||||
|
||||
if (type != NULL_TREE)
|
||||
m = mangle_type_string (TYPE_MAIN_VARIANT (type));
|
||||
|
||||
const char prefix[] = "omp declare mapper ";
|
||||
size_t lenp = sizeof (prefix);
|
||||
if (strncmp (p, prefix, lenp - 1) == 0)
|
||||
lenp = 1;
|
||||
size_t len = strlen (p);
|
||||
size_t lenm = m ? strlen (m) + 1 : 0;
|
||||
char *name = XALLOCAVEC (char, lenp + len + lenm);
|
||||
memcpy (name, prefix, lenp - 1);
|
||||
memcpy (name + lenp - 1, p, len + 1);
|
||||
if (m)
|
||||
{
|
||||
name[lenp + len - 1] = '~';
|
||||
memcpy (name + lenp + len, m, lenm);
|
||||
}
|
||||
return get_identifier (name);
|
||||
}
|
||||
|
||||
tree
|
||||
cxx_omp_mapper_lookup (tree id, tree type)
|
||||
{
|
||||
if (!RECORD_OR_UNION_TYPE_P (type))
|
||||
return NULL_TREE;
|
||||
id = omp_mapper_id (id, type);
|
||||
return lookup_name (id);
|
||||
}
|
||||
|
||||
tree
|
||||
cxx_omp_extract_mapper_directive (tree vardecl)
|
||||
{
|
||||
gcc_assert (TREE_CODE (vardecl) == VAR_DECL);
|
||||
|
||||
/* Instantiate the decl if we haven't already. */
|
||||
mark_used (vardecl);
|
||||
tree body = DECL_INITIAL (vardecl);
|
||||
|
||||
if (TREE_CODE (body) == STATEMENT_LIST)
|
||||
{
|
||||
tree_stmt_iterator tsi = tsi_start (body);
|
||||
gcc_assert (TREE_CODE (tsi_stmt (tsi)) == DECL_EXPR);
|
||||
tsi_next (&tsi);
|
||||
body = tsi_stmt (tsi);
|
||||
}
|
||||
|
||||
gcc_assert (TREE_CODE (body) == OMP_DECLARE_MAPPER);
|
||||
|
||||
return body;
|
||||
}
|
||||
|
||||
/* For now we can handle singleton OMP_ARRAY_SECTIONs with custom mappers, but
|
||||
nothing more complicated. */
|
||||
|
||||
tree
|
||||
cxx_omp_map_array_section (location_t loc, tree t)
|
||||
{
|
||||
tree low = TREE_OPERAND (t, 1);
|
||||
tree len = TREE_OPERAND (t, 2);
|
||||
|
||||
if (len && integer_onep (len))
|
||||
{
|
||||
t = TREE_OPERAND (t, 0);
|
||||
|
||||
if (!low)
|
||||
low = integer_zero_node;
|
||||
|
||||
if (TREE_CODE (TREE_TYPE (t)) == REFERENCE_TYPE)
|
||||
t = convert_from_reference (t);
|
||||
|
||||
t = build_array_ref (loc, t, low);
|
||||
}
|
||||
|
||||
return t;
|
||||
}
|
||||
|
||||
/* Helper function for cp_parser_omp_declare_reduction_exprs
|
||||
and tsubst_omp_udr.
|
||||
Remove CLEANUP_STMT for data (omp_priv variable).
|
||||
|
|
@ -7221,6 +7313,33 @@ finish_omp_reduction_clause (tree c, bool *need_default_ctor, bool *need_dtor)
|
|||
return false;
|
||||
}
|
||||
|
||||
/* Check an instance of an "omp declare mapper" function. */
|
||||
|
||||
bool
|
||||
cp_check_omp_declare_mapper (tree udm)
|
||||
{
|
||||
tree type = TREE_TYPE (udm);
|
||||
location_t loc = DECL_SOURCE_LOCATION (udm);
|
||||
|
||||
if (type == error_mark_node)
|
||||
return false;
|
||||
|
||||
if (!processing_template_decl && !RECORD_OR_UNION_TYPE_P (type))
|
||||
{
|
||||
error_at (loc, "%qT is not a struct, union or class type in "
|
||||
"%<#pragma omp declare mapper%>", type);
|
||||
return false;
|
||||
}
|
||||
if (!processing_template_decl && CLASSTYPE_VBASECLASSES (type))
|
||||
{
|
||||
error_at (loc, "%qT must not be a virtual base class in "
|
||||
"%<#pragma omp declare mapper%>", type);
|
||||
return false;
|
||||
}
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
/* Called from finish_struct_1. linear(this) or linear(this:step)
|
||||
clauses might not be finalized yet because the class has been incomplete
|
||||
when parsing #pragma omp declare simd methods. Fix those up now. */
|
||||
|
|
@ -8865,6 +8984,12 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
|
|||
case OMP_CLAUSE_MAP:
|
||||
if (OMP_CLAUSE_MAP_IMPLICIT (c) && !implicit_moved)
|
||||
goto move_implicit;
|
||||
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME
|
||||
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POP_MAPPER_NAME)
|
||||
{
|
||||
remove = true;
|
||||
break;
|
||||
}
|
||||
/* FALLTHRU */
|
||||
case OMP_CLAUSE_TO:
|
||||
case OMP_CLAUSE_FROM:
|
||||
|
|
@ -10484,6 +10609,8 @@ struct omp_target_walk_data
|
|||
/* Local variables declared inside a BIND_EXPR, used to filter out such
|
||||
variables when recording lambda_objects_accessed. */
|
||||
hash_set<tree> local_decls;
|
||||
|
||||
omp_mapper_list<tree> *mappers;
|
||||
};
|
||||
|
||||
/* Helper function of finish_omp_target_clauses, called via
|
||||
|
|
@ -10497,6 +10624,7 @@ finish_omp_target_clauses_r (tree *tp, int *walk_subtrees, void *ptr)
|
|||
struct omp_target_walk_data *data = (struct omp_target_walk_data *) ptr;
|
||||
tree current_object = data->current_object;
|
||||
tree current_closure = data->current_closure;
|
||||
omp_mapper_list<tree> *mlist = data->mappers;
|
||||
|
||||
/* References inside of these expression codes shouldn't incur any
|
||||
form of mapping, so return early. */
|
||||
|
|
@ -10510,6 +10638,27 @@ finish_omp_target_clauses_r (tree *tp, int *walk_subtrees, void *ptr)
|
|||
if (TREE_CODE (t) == OMP_CLAUSE)
|
||||
return NULL_TREE;
|
||||
|
||||
if (!processing_template_decl)
|
||||
{
|
||||
tree aggr_type = NULL_TREE;
|
||||
|
||||
if (TREE_CODE (t) == COMPONENT_REF
|
||||
&& RECORD_OR_UNION_TYPE_P (TREE_TYPE (TREE_OPERAND (t, 0))))
|
||||
aggr_type = TREE_TYPE (TREE_OPERAND (t, 0));
|
||||
else if ((TREE_CODE (t) == VAR_DECL
|
||||
|| TREE_CODE (t) == PARM_DECL
|
||||
|| TREE_CODE (t) == RESULT_DECL)
|
||||
&& RECORD_OR_UNION_TYPE_P (TREE_TYPE (t)))
|
||||
aggr_type = TREE_TYPE (t);
|
||||
|
||||
if (aggr_type)
|
||||
{
|
||||
tree mapper_fn = cxx_omp_mapper_lookup (NULL_TREE, aggr_type);
|
||||
if (mapper_fn)
|
||||
mlist->add_mapper (NULL_TREE, aggr_type, mapper_fn);
|
||||
}
|
||||
}
|
||||
|
||||
if (current_object)
|
||||
{
|
||||
tree this_expr = TREE_OPERAND (current_object, 0);
|
||||
|
|
@ -10612,10 +10761,48 @@ finish_omp_target_clauses (location_t loc, tree body, tree *clauses_ptr)
|
|||
else
|
||||
data.current_closure = NULL_TREE;
|
||||
|
||||
cp_walk_tree_without_duplicates (&body, finish_omp_target_clauses_r, &data);
|
||||
|
||||
auto_vec<tree, 16> new_clauses;
|
||||
|
||||
if (!processing_template_decl)
|
||||
{
|
||||
hash_set<omp_name_type<tree> > seen_types;
|
||||
auto_vec<tree> mapper_fns;
|
||||
omp_mapper_list<tree> mlist (&seen_types, &mapper_fns);
|
||||
data.mappers = &mlist;
|
||||
|
||||
cp_walk_tree_without_duplicates (&body, finish_omp_target_clauses_r,
|
||||
&data);
|
||||
|
||||
unsigned int i;
|
||||
tree mapper_fn;
|
||||
FOR_EACH_VEC_ELT (mapper_fns, i, mapper_fn)
|
||||
c_omp_find_nested_mappers (&mlist, mapper_fn);
|
||||
|
||||
FOR_EACH_VEC_ELT (mapper_fns, i, mapper_fn)
|
||||
{
|
||||
tree mapper = cxx_omp_extract_mapper_directive (mapper_fn);
|
||||
if (mapper == error_mark_node)
|
||||
continue;
|
||||
tree mapper_name = OMP_DECLARE_MAPPER_ID (mapper);
|
||||
tree decl = OMP_DECLARE_MAPPER_DECL (mapper);
|
||||
if (BASELINK_P (mapper_fn))
|
||||
mapper_fn = BASELINK_FUNCTIONS (mapper_fn);
|
||||
|
||||
tree c = build_omp_clause (loc, OMP_CLAUSE__MAPPER_BINDING_);
|
||||
OMP_CLAUSE__MAPPER_BINDING__ID (c) = mapper_name;
|
||||
OMP_CLAUSE__MAPPER_BINDING__DECL (c) = decl;
|
||||
OMP_CLAUSE__MAPPER_BINDING__MAPPER (c) = mapper_fn;
|
||||
|
||||
new_clauses.safe_push (c);
|
||||
}
|
||||
}
|
||||
else
|
||||
{
|
||||
data.mappers = NULL;
|
||||
cp_walk_tree_without_duplicates (&body, finish_omp_target_clauses_r,
|
||||
&data);
|
||||
}
|
||||
|
||||
tree omp_target_this_expr = NULL_TREE;
|
||||
tree *explicit_this_deref_map = NULL;
|
||||
if (data.this_expr_accessed)
|
||||
|
|
|
|||
|
|
@ -27,6 +27,9 @@ along with GCC; see the file COPYING3. If not see
|
|||
#include "match.h"
|
||||
#include "parse.h"
|
||||
#include "tree-core.h"
|
||||
#include "tree.h"
|
||||
#include "fold-const.h"
|
||||
#include "tree-hash-traits.h"
|
||||
#include "omp-general.h"
|
||||
|
||||
/* Current statement label. Zero means no statement label. Because new_st
|
||||
|
|
|
|||
263
gcc/gimplify.cc
263
gcc/gimplify.cc
|
|
@ -273,6 +273,7 @@ struct gimplify_omp_ctx
|
|||
{
|
||||
struct gimplify_omp_ctx *outer_context;
|
||||
splay_tree variables;
|
||||
hash_map<omp_name_type<tree>, tree> *implicit_mappers;
|
||||
hash_set<tree> *privatized_types;
|
||||
tree clauses;
|
||||
/* Iteration variables in an OMP_FOR. */
|
||||
|
|
@ -507,6 +508,7 @@ new_omp_context (enum omp_region_type region_type)
|
|||
c = XCNEW (struct gimplify_omp_ctx);
|
||||
c->outer_context = gimplify_omp_ctxp;
|
||||
c->variables = splay_tree_new (splay_tree_compare_decl_uid, 0, 0);
|
||||
c->implicit_mappers = new hash_map<omp_name_type<tree>, tree>;
|
||||
c->privatized_types = new hash_set<tree>;
|
||||
c->location = input_location;
|
||||
c->region_type = region_type;
|
||||
|
|
@ -530,6 +532,7 @@ delete_omp_context (struct gimplify_omp_ctx *c)
|
|||
{
|
||||
splay_tree_delete (c->variables);
|
||||
delete c->privatized_types;
|
||||
delete c->implicit_mappers;
|
||||
c->loop_iter_var.release ();
|
||||
XDELETE (c);
|
||||
}
|
||||
|
|
@ -12892,6 +12895,218 @@ error_out:
|
|||
return success;
|
||||
}
|
||||
|
||||
struct instantiate_mapper_info
|
||||
{
|
||||
tree *mapper_clauses_p;
|
||||
struct gimplify_omp_ctx *omp_ctx;
|
||||
gimple_seq *pre_p;
|
||||
};
|
||||
|
||||
/* Helper function for omp_instantiate_mapper. */
|
||||
|
||||
static tree
|
||||
remap_mapper_decl_1 (tree *tp, int *walk_subtrees, void *data)
|
||||
{
|
||||
copy_body_data *id = (copy_body_data *) data;
|
||||
|
||||
if (DECL_P (*tp))
|
||||
{
|
||||
tree replacement = remap_decl (*tp, id);
|
||||
if (*tp != replacement)
|
||||
{
|
||||
*tp = unshare_expr (replacement);
|
||||
*walk_subtrees = 0;
|
||||
}
|
||||
}
|
||||
|
||||
return NULL_TREE;
|
||||
}
|
||||
|
||||
/* A copy_decl implementation (for use with tree-inline.cc functions) that
|
||||
only transform decls or SSA names that are part of a map we already
|
||||
prepared. */
|
||||
|
||||
static tree
|
||||
omp_mapper_copy_decl (tree var, copy_body_data *cb)
|
||||
{
|
||||
tree *repl = cb->decl_map->get (var);
|
||||
|
||||
if (repl)
|
||||
return *repl;
|
||||
|
||||
return var;
|
||||
}
|
||||
|
||||
static tree *
|
||||
omp_instantiate_mapper (gimple_seq *pre_p,
|
||||
hash_map<omp_name_type<tree>, tree> *implicit_mappers,
|
||||
tree mapperfn, tree expr, enum gomp_map_kind outer_kind,
|
||||
tree *mapper_clauses_p)
|
||||
{
|
||||
tree mapper_name = NULL_TREE;
|
||||
tree mapper = lang_hooks.decls.omp_extract_mapper_directive (mapperfn);
|
||||
gcc_assert (TREE_CODE (mapper) == OMP_DECLARE_MAPPER);
|
||||
|
||||
tree clause = OMP_DECLARE_MAPPER_CLAUSES (mapper);
|
||||
tree dummy_var = OMP_DECLARE_MAPPER_DECL (mapper);
|
||||
|
||||
/* The "extraction map" is used to map the mapper variable in the "declare
|
||||
mapper" directive, and also any temporary variables that have been created
|
||||
as part of expanding the mapper function's body (which are expanded as a
|
||||
"bind" expression in the pre_p sequence). */
|
||||
hash_map<tree, tree> extraction_map;
|
||||
|
||||
extraction_map.put (dummy_var, expr);
|
||||
extraction_map.put (expr, expr);
|
||||
|
||||
/* This copy_body_data is only used to remap the decls in the
|
||||
OMP_DECLARE_MAPPER tree node expansion itself. All relevant decls should
|
||||
already be in the current function. */
|
||||
copy_body_data id;
|
||||
memset (&id, 0, sizeof (id));
|
||||
id.src_fn = current_function_decl;
|
||||
id.dst_fn = current_function_decl;
|
||||
id.src_cfun = cfun;
|
||||
id.decl_map = &extraction_map;
|
||||
id.copy_decl = omp_mapper_copy_decl;
|
||||
id.transform_call_graph_edges = CB_CGE_DUPLICATE; // ???
|
||||
id.transform_new_cfg = true; // ???
|
||||
|
||||
for (; clause; clause = OMP_CLAUSE_CHAIN (clause))
|
||||
{
|
||||
enum gomp_map_kind map_kind = OMP_CLAUSE_MAP_KIND (clause);
|
||||
tree *nested_mapper_p = NULL;
|
||||
|
||||
if (map_kind == GOMP_MAP_PUSH_MAPPER_NAME)
|
||||
{
|
||||
mapper_name = OMP_CLAUSE_DECL (clause);
|
||||
continue;
|
||||
}
|
||||
else if (map_kind == GOMP_MAP_POP_MAPPER_NAME)
|
||||
{
|
||||
mapper_name = NULL_TREE;
|
||||
continue;
|
||||
}
|
||||
|
||||
tree decl = OMP_CLAUSE_DECL (clause);
|
||||
tree unshared, type;
|
||||
bool nonunit_array_with_mapper = false;
|
||||
|
||||
if (TREE_CODE (decl) == OMP_ARRAY_SECTION)
|
||||
{
|
||||
location_t loc = OMP_CLAUSE_LOCATION (clause);
|
||||
tree tmp = lang_hooks.decls.omp_map_array_section (loc, decl);
|
||||
if (tmp == decl)
|
||||
{
|
||||
unshared = unshare_expr (clause);
|
||||
nonunit_array_with_mapper = true;
|
||||
type = TREE_TYPE (TREE_TYPE (decl));
|
||||
}
|
||||
else
|
||||
{
|
||||
unshared = build_omp_clause (OMP_CLAUSE_LOCATION (clause),
|
||||
OMP_CLAUSE_CODE (clause));
|
||||
OMP_CLAUSE_DECL (unshared) = tmp;
|
||||
OMP_CLAUSE_SIZE (unshared)
|
||||
= DECL_P (tmp) ? DECL_SIZE_UNIT (tmp)
|
||||
: TYPE_SIZE_UNIT (TREE_TYPE (tmp));
|
||||
type = TREE_TYPE (tmp);
|
||||
}
|
||||
}
|
||||
else
|
||||
{
|
||||
unshared = unshare_expr (clause);
|
||||
type = TREE_TYPE (decl);
|
||||
}
|
||||
|
||||
walk_tree (&unshared, remap_mapper_decl_1, &id, NULL);
|
||||
|
||||
if (OMP_CLAUSE_MAP_KIND (unshared) == GOMP_MAP_UNSET)
|
||||
OMP_CLAUSE_SET_MAP_KIND (unshared, outer_kind);
|
||||
|
||||
decl = OMP_CLAUSE_DECL (unshared);
|
||||
type = TYPE_MAIN_VARIANT (type);
|
||||
|
||||
nested_mapper_p = implicit_mappers->get ({ mapper_name, type });
|
||||
|
||||
if (nested_mapper_p && *nested_mapper_p != mapperfn)
|
||||
{
|
||||
if (nonunit_array_with_mapper)
|
||||
{
|
||||
sorry ("user-defined mapper with non-unit length array section");
|
||||
continue;
|
||||
}
|
||||
|
||||
if (map_kind == GOMP_MAP_UNSET)
|
||||
map_kind = outer_kind;
|
||||
|
||||
mapper_clauses_p
|
||||
= omp_instantiate_mapper (pre_p, implicit_mappers,
|
||||
*nested_mapper_p, decl, map_kind,
|
||||
mapper_clauses_p);
|
||||
continue;
|
||||
}
|
||||
|
||||
*mapper_clauses_p = unshared;
|
||||
mapper_clauses_p = &OMP_CLAUSE_CHAIN (unshared);
|
||||
}
|
||||
|
||||
return mapper_clauses_p;
|
||||
}
|
||||
|
||||
static int
|
||||
omp_instantiate_implicit_mappers (splay_tree_node n, void *data)
|
||||
{
|
||||
tree decl = (tree) n->key;
|
||||
instantiate_mapper_info *im_info = (instantiate_mapper_info *) data;
|
||||
gimplify_omp_ctx *ctx = im_info->omp_ctx;
|
||||
tree *mapper_p = NULL;
|
||||
tree type = TREE_TYPE (decl);
|
||||
bool ref_p = false;
|
||||
unsigned flags = n->value;
|
||||
|
||||
if (flags & (GOVD_EXPLICIT | GOVD_LOCAL))
|
||||
return 0;
|
||||
if ((flags & GOVD_SEEN) == 0)
|
||||
return 0;
|
||||
/* If we already have clauses pertaining to a struct variable, then we don't
|
||||
want to implicitly invoke a user-defined mapper. */
|
||||
if ((flags & GOVD_EXPLICIT) != 0 && AGGREGATE_TYPE_P (TREE_TYPE (decl)))
|
||||
return 0;
|
||||
|
||||
if (TREE_CODE (type) == REFERENCE_TYPE)
|
||||
{
|
||||
ref_p = true;
|
||||
type = TREE_TYPE (type);
|
||||
}
|
||||
|
||||
type = TYPE_MAIN_VARIANT (type);
|
||||
|
||||
if (DECL_P (decl) && type && AGGREGATE_TYPE_P (type))
|
||||
{
|
||||
gcc_assert (ctx);
|
||||
mapper_p = ctx->implicit_mappers->get ({ NULL_TREE, type });
|
||||
}
|
||||
|
||||
if (mapper_p)
|
||||
{
|
||||
/* If we have a reference, map the pointed-to object rather than the
|
||||
reference itself. */
|
||||
if (ref_p)
|
||||
decl = build_fold_indirect_ref (decl);
|
||||
|
||||
im_info->mapper_clauses_p
|
||||
= omp_instantiate_mapper (im_info->pre_p, ctx->implicit_mappers,
|
||||
*mapper_p, decl, GOMP_MAP_TOFROM,
|
||||
im_info->mapper_clauses_p);
|
||||
/* Make sure we don't map the same variable implicitly in
|
||||
gimplify_adjust_omp_clauses_1 also. */
|
||||
n->value |= GOVD_EXPLICIT;
|
||||
}
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
/* Scan the OMP clauses in *LIST_P, installing mappings into a new
|
||||
and previous omp contexts. */
|
||||
|
||||
|
|
@ -13688,6 +13903,17 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
|
|||
}
|
||||
goto do_notice;
|
||||
|
||||
case OMP_CLAUSE__MAPPER_BINDING_:
|
||||
{
|
||||
tree name = OMP_CLAUSE__MAPPER_BINDING__ID (c);
|
||||
tree var = OMP_CLAUSE__MAPPER_BINDING__DECL (c);
|
||||
tree type = TYPE_MAIN_VARIANT (TREE_TYPE (var));
|
||||
tree fndecl = OMP_CLAUSE__MAPPER_BINDING__MAPPER (c);
|
||||
ctx->implicit_mappers->put ({ name, type }, fndecl);
|
||||
remove = true;
|
||||
break;
|
||||
}
|
||||
|
||||
case OMP_CLAUSE_USE_DEVICE_PTR:
|
||||
case OMP_CLAUSE_USE_DEVICE_ADDR:
|
||||
flags = GOVD_EXPLICIT;
|
||||
|
|
@ -14772,6 +14998,30 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
|
|||
|| code == OMP_TARGET_ENTER_DATA
|
||||
|| code == OMP_TARGET_EXIT_DATA)
|
||||
{
|
||||
tree mapper_clauses = NULL_TREE;
|
||||
instantiate_mapper_info im_info;
|
||||
|
||||
im_info.mapper_clauses_p = &mapper_clauses;
|
||||
im_info.omp_ctx = ctx;
|
||||
im_info.pre_p = pre_p;
|
||||
|
||||
splay_tree_foreach (ctx->variables,
|
||||
omp_instantiate_implicit_mappers,
|
||||
(void *) &im_info);
|
||||
|
||||
if (mapper_clauses)
|
||||
{
|
||||
mapper_clauses
|
||||
= lang_hooks.decls.omp_finish_mapper_clauses (mapper_clauses);
|
||||
|
||||
/* Stick the implicitly-expanded mapper clauses at the end of the
|
||||
clause list. */
|
||||
tree *tail = list_p;
|
||||
while (*tail)
|
||||
tail = &OMP_CLAUSE_CHAIN (*tail);
|
||||
*tail = mapper_clauses;
|
||||
}
|
||||
|
||||
vec<omp_mapping_group> *groups;
|
||||
groups = omp_gather_mapping_groups (list_p);
|
||||
hash_map<tree_operand_hash_no_se, omp_mapping_group *> *grpmap = NULL;
|
||||
|
|
@ -19257,6 +19507,15 @@ gimplify_omp_metadirective (tree *expr_p, gimple_seq *pre_p, gimple_seq *,
|
|||
return GS_OK;
|
||||
}
|
||||
|
||||
/* Gimplify an OMP_DECLARE_MAPPER node (by just removing it). */
|
||||
|
||||
static enum gimplify_status
|
||||
gimplify_omp_declare_mapper (tree *expr_p)
|
||||
{
|
||||
*expr_p = NULL_TREE;
|
||||
return GS_ALL_DONE;
|
||||
}
|
||||
|
||||
/* Convert the GENERIC expression tree *EXPR_P to GIMPLE. If the
|
||||
expression produces a value to be used as an operand inside a GIMPLE
|
||||
statement, the value will be stored back in *EXPR_P. This value will
|
||||
|
|
@ -20218,6 +20477,10 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
|
|||
ret = GS_ALL_DONE;
|
||||
break;
|
||||
|
||||
case OMP_DECLARE_MAPPER:
|
||||
ret = gimplify_omp_declare_mapper (expr_p);
|
||||
break;
|
||||
|
||||
case TRANSACTION_EXPR:
|
||||
ret = gimplify_transaction (expr_p, pre_p);
|
||||
break;
|
||||
|
|
|
|||
|
|
@ -90,6 +90,9 @@ extern bool lhd_omp_deep_mapping_p (const gimple *, tree);
|
|||
extern tree lhd_omp_deep_mapping_cnt (const gimple *, tree, gimple_seq *);
|
||||
extern void lhd_omp_deep_mapping (const gimple *, tree, unsigned HOST_WIDE_INT,
|
||||
tree, tree, tree, tree, tree, gimple_seq *);
|
||||
extern tree lhd_omp_mapper_lookup (tree, tree);
|
||||
extern tree lhd_omp_extract_mapper_directive (tree);
|
||||
extern tree lhd_omp_map_array_section (location_t, tree);
|
||||
struct gimplify_omp_ctx;
|
||||
extern void lhd_omp_firstprivatize_type_sizes (struct gimplify_omp_ctx *,
|
||||
tree);
|
||||
|
|
@ -279,6 +282,11 @@ extern tree lhd_unit_size_without_reusable_padding (tree);
|
|||
#define LANG_HOOKS_OMP_DEEP_MAPPING_P lhd_omp_deep_mapping_p
|
||||
#define LANG_HOOKS_OMP_DEEP_MAPPING_CNT lhd_omp_deep_mapping_cnt
|
||||
#define LANG_HOOKS_OMP_DEEP_MAPPING lhd_omp_deep_mapping
|
||||
#define LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES lhd_pass_through_t
|
||||
#define LANG_HOOKS_OMP_MAPPER_LOOKUP lhd_omp_mapper_lookup
|
||||
#define LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE \
|
||||
lhd_omp_extract_mapper_directive
|
||||
#define LANG_HOOKS_OMP_MAP_ARRAY_SECTION lhd_omp_map_array_section
|
||||
#define LANG_HOOKS_OMP_ALLOCATABLE_P hook_bool_tree_false
|
||||
#define LANG_HOOKS_OMP_SCALAR_P lhd_omp_scalar_p
|
||||
#define LANG_HOOKS_OMP_SCALAR_TARGET_P hook_bool_tree_false
|
||||
|
|
@ -316,6 +324,10 @@ extern tree lhd_unit_size_without_reusable_padding (tree);
|
|||
LANG_HOOKS_OMP_DEEP_MAPPING_P, \
|
||||
LANG_HOOKS_OMP_DEEP_MAPPING_CNT, \
|
||||
LANG_HOOKS_OMP_DEEP_MAPPING, \
|
||||
LANG_HOOKS_OMP_FINISH_MAPPER_CLAUSES, \
|
||||
LANG_HOOKS_OMP_MAPPER_LOOKUP, \
|
||||
LANG_HOOKS_OMP_EXTRACT_MAPPER_DIRECTIVE, \
|
||||
LANG_HOOKS_OMP_MAP_ARRAY_SECTION, \
|
||||
LANG_HOOKS_OMP_ALLOCATABLE_P, \
|
||||
LANG_HOOKS_OMP_SCALAR_P, \
|
||||
LANG_HOOKS_OMP_SCALAR_TARGET_P, \
|
||||
|
|
|
|||
|
|
@ -669,6 +669,32 @@ lhd_omp_deep_mapping (const gimple *, tree, unsigned HOST_WIDE_INT, tree, tree,
|
|||
{
|
||||
}
|
||||
|
||||
/* Look up an OpenMP "declare mapper" mapper. */
|
||||
|
||||
tree
|
||||
lhd_omp_mapper_lookup (tree, tree)
|
||||
{
|
||||
return NULL_TREE;
|
||||
}
|
||||
|
||||
/* Given the representation used by the front-end to contain a mapper
|
||||
directive, return the statement for the directive itself. */
|
||||
|
||||
tree
|
||||
lhd_omp_extract_mapper_directive (tree)
|
||||
{
|
||||
return error_mark_node;
|
||||
}
|
||||
|
||||
/* Return a simplified form for OMP_ARRAY_SECTION argument, or
|
||||
error_mark_node if impossible. */
|
||||
|
||||
tree
|
||||
lhd_omp_map_array_section (location_t, tree)
|
||||
{
|
||||
return error_mark_node;
|
||||
}
|
||||
|
||||
/* Return true if DECL is a scalar variable (for the purpose of
|
||||
implicit firstprivatization & mapping). Only if alloc_ptr_ok
|
||||
are allocatables and pointers accepted. */
|
||||
|
|
|
|||
|
|
@ -328,6 +328,22 @@ struct lang_hooks_for_decls
|
|||
tree data, tree sizes, tree kinds,
|
||||
tree offset_data, tree offset, gimple_seq *seq);
|
||||
|
||||
/* Finish language-specific processing on mapping nodes after expanding
|
||||
user-defined mappers. */
|
||||
tree (*omp_finish_mapper_clauses) (tree clauses);
|
||||
|
||||
/* Find a mapper in the current parsing context, given a NAME (or
|
||||
NULL_TREE) and TYPE. */
|
||||
tree (*omp_mapper_lookup) (tree name, tree type);
|
||||
|
||||
/* Return the statement for the mapper directive definition, from the
|
||||
representation used to contain it (e.g. an inline function
|
||||
declaration). */
|
||||
tree (*omp_extract_mapper_directive) (tree fndecl);
|
||||
|
||||
/* Return a simplified form for OMP_ARRAY_SECTION argument. */
|
||||
tree (*omp_map_array_section) (location_t, tree t);
|
||||
|
||||
/* Return true if DECL is an allocatable variable (for the purpose of
|
||||
implicit mapping). */
|
||||
bool (*omp_allocatable_p) (tree decl);
|
||||
|
|
|
|||
|
|
@ -254,6 +254,92 @@ get_openacc_privatization_dump_flags ()
|
|||
|
||||
extern tree omp_build_component_ref (tree obj, tree field);
|
||||
|
||||
template <typename T>
|
||||
struct omp_name_type
|
||||
{
|
||||
tree name;
|
||||
T type;
|
||||
};
|
||||
|
||||
template <>
|
||||
struct default_hash_traits <omp_name_type<tree> >
|
||||
: typed_noop_remove <omp_name_type<tree> >
|
||||
{
|
||||
GTY((skip)) typedef omp_name_type<tree> value_type;
|
||||
GTY((skip)) typedef omp_name_type<tree> compare_type;
|
||||
|
||||
static hashval_t
|
||||
hash (omp_name_type<tree> p)
|
||||
{
|
||||
return p.name ? iterative_hash_expr (p.name, TYPE_UID (p.type))
|
||||
: TYPE_UID (p.type);
|
||||
}
|
||||
|
||||
static const bool empty_zero_p = true;
|
||||
|
||||
static bool
|
||||
is_empty (omp_name_type<tree> p)
|
||||
{
|
||||
return p.type == NULL;
|
||||
}
|
||||
|
||||
static bool
|
||||
is_deleted (omp_name_type<tree>)
|
||||
{
|
||||
return false;
|
||||
}
|
||||
|
||||
static bool
|
||||
equal (const omp_name_type<tree> &a, const omp_name_type<tree> &b)
|
||||
{
|
||||
if (a.name == NULL_TREE && b.name == NULL_TREE)
|
||||
return a.type == b.type;
|
||||
else if (a.name == NULL_TREE || b.name == NULL_TREE)
|
||||
return false;
|
||||
else
|
||||
return a.name == b.name && a.type == b.type;
|
||||
}
|
||||
|
||||
static void
|
||||
mark_empty (omp_name_type<tree> &e)
|
||||
{
|
||||
e.type = NULL;
|
||||
}
|
||||
};
|
||||
|
||||
template <typename T>
|
||||
struct omp_mapper_list
|
||||
{
|
||||
hash_set<omp_name_type<T>> *seen_types;
|
||||
vec<tree> *mappers;
|
||||
|
||||
omp_mapper_list (hash_set<omp_name_type<T>> *s, vec<tree> *m)
|
||||
: seen_types (s), mappers (m) { }
|
||||
|
||||
void add_mapper (tree name, T type, tree mapperfn)
|
||||
{
|
||||
/* We can't hash a NULL_TREE... */
|
||||
if (!name)
|
||||
name = void_node;
|
||||
|
||||
omp_name_type<T> n_t = { name, type };
|
||||
|
||||
if (seen_types->contains (n_t))
|
||||
return;
|
||||
|
||||
seen_types->add (n_t);
|
||||
mappers->safe_push (mapperfn);
|
||||
}
|
||||
|
||||
bool contains (tree name, T type)
|
||||
{
|
||||
if (!name)
|
||||
name = void_node;
|
||||
|
||||
return seen_types->contains ({ name, type });
|
||||
}
|
||||
};
|
||||
|
||||
namespace omp_addr_tokenizer {
|
||||
|
||||
/* These are the ways of accessing a variable that have special-case handling
|
||||
|
|
|
|||
|
|
@ -0,0 +1,30 @@
|
|||
/* { dg-do compile { target c++ } } */
|
||||
/* { dg-additional-options "-fdump-tree-gimple" } */
|
||||
|
||||
struct S {
|
||||
int x,y;
|
||||
};
|
||||
|
||||
#pragma omp declare mapper(default : struct S var) map(mapper(default), tofrom: var)
|
||||
#pragma omp declare mapper(only_x : struct S var) map(mapper(default), tofrom: var.x)
|
||||
|
||||
void f(){
|
||||
struct S z = {1,2};
|
||||
#pragma omp target defaultmap(alloc)
|
||||
z.x += 5;
|
||||
#pragma omp target map(z)
|
||||
z.x += 7;
|
||||
#pragma omp target map(mapper(default), tofrom: z)
|
||||
z.x += 8;
|
||||
#pragma omp target map(mapper(only_x), tofrom: z)
|
||||
z.x += 9;
|
||||
if (z.x != 1+5+7+8+9) __builtin_abort ();
|
||||
}
|
||||
|
||||
/* { dg-final { scan-tree-dump-times "#pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) defaultmap\\(alloc\\) map\\(tofrom:z \\\[len: 8\\\]\\)" 1 "gimple" } } */
|
||||
/* { dg-final { scan-tree-dump-times "#pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) map\\(tofrom:z \\\[len: 8\\\]\\)" 2 "gimple" } } */
|
||||
/* { dg-final { scan-tree-dump-times "#pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) map\\(struct:z \\\[len: 1\\\]\\) map\\(tofrom:z.x \\\[len: 4\\\]\\)" 1 "gimple" } } */
|
||||
|
||||
int main() {
|
||||
f();
|
||||
}
|
||||
|
|
@ -0,0 +1,22 @@
|
|||
/* { dg-do compile { target c++ } } */
|
||||
|
||||
struct XYZ {
|
||||
int a;
|
||||
int *b;
|
||||
int c;
|
||||
};
|
||||
|
||||
#pragma omp declare mapper(struct XYZ t)
|
||||
/* { dg-error "missing 'map' clause" "" { target c } .-1 } */
|
||||
/* { dg-error "missing 'map' clause before end of line" "" { target c++ } .-2 } */
|
||||
|
||||
struct ABC {
|
||||
int *a;
|
||||
int b;
|
||||
int c;
|
||||
};
|
||||
|
||||
#pragma omp declare mapper(struct ABC d) firstprivate(d.b)
|
||||
/* { dg-error "unexpected clause" "" { target c } .-1 } */
|
||||
/* { dg-error "expected end of line before '\\(' token" "" { target c } .-2 } */
|
||||
/* { dg-error "unexpected clause before '\\(' token" "" { target c++ } .-3 } */
|
||||
|
|
@ -0,0 +1,30 @@
|
|||
// { dg-do compile { target c++ } }
|
||||
// { dg-additional-options "-fdump-tree-gimple" }
|
||||
|
||||
#include <stdlib.h>
|
||||
|
||||
// Test named mapper invocation.
|
||||
|
||||
struct S {
|
||||
int *ptr;
|
||||
int size;
|
||||
};
|
||||
|
||||
int main (int argc, char *argv[])
|
||||
{
|
||||
int N = 1024;
|
||||
#pragma omp declare mapper (mapN:struct S s) map(to:s.ptr, s.size) \
|
||||
map(s.ptr[:N])
|
||||
|
||||
struct S s;
|
||||
s.ptr = (int *) malloc (sizeof (int) * N);
|
||||
|
||||
#pragma omp target map(mapper(mapN), tofrom: s)
|
||||
// { dg-final { scan-tree-dump {map\(struct:s \[len: 2\]\) map\(alloc:s\.ptr \[len: [0-9]+\]\) map\(to:s\.size \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:s\.ptr \[bias: 0\]\)} "gimple" } }
|
||||
{
|
||||
for (int i = 0; i < N; i++)
|
||||
s.ptr[i]++;
|
||||
}
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
|
@ -0,0 +1,78 @@
|
|||
/* { dg-do compile { target c++ } } */
|
||||
/* { dg-additional-options "-fdump-tree-original" } */
|
||||
|
||||
/* Check mapper binding clauses. */
|
||||
|
||||
struct Y {
|
||||
int z;
|
||||
};
|
||||
|
||||
struct Z {
|
||||
int z;
|
||||
};
|
||||
|
||||
#pragma omp declare mapper (struct Y y) map(tofrom: y)
|
||||
#pragma omp declare mapper (struct Z z) map(tofrom: z)
|
||||
|
||||
int foo (void)
|
||||
{
|
||||
struct Y yy;
|
||||
struct Z zz;
|
||||
int dummy;
|
||||
|
||||
#pragma omp target data map(dummy)
|
||||
{
|
||||
#pragma omp target
|
||||
{
|
||||
yy.z++;
|
||||
zz.z++;
|
||||
}
|
||||
yy.z++;
|
||||
}
|
||||
return yy.z;
|
||||
}
|
||||
|
||||
struct P
|
||||
{
|
||||
struct Z *zp;
|
||||
};
|
||||
|
||||
int bar (void)
|
||||
{
|
||||
struct Y yy;
|
||||
struct Z zz;
|
||||
struct P pp;
|
||||
struct Z t;
|
||||
int dummy;
|
||||
|
||||
pp.zp = &t;
|
||||
|
||||
#pragma omp declare mapper (struct Y y) map(tofrom: y.z)
|
||||
#pragma omp declare mapper (struct Z z) map(tofrom: z.z)
|
||||
|
||||
#pragma omp target data map(dummy)
|
||||
{
|
||||
#pragma omp target
|
||||
{
|
||||
yy.z++;
|
||||
zz.z++;
|
||||
}
|
||||
yy.z++;
|
||||
}
|
||||
|
||||
#pragma omp declare mapper(struct P x) map(to:x.zp) map(tofrom:*x.zp)
|
||||
|
||||
#pragma omp target
|
||||
{
|
||||
zz = *pp.zp;
|
||||
}
|
||||
|
||||
return zz.z;
|
||||
}
|
||||
|
||||
/* { dg-final { scan-tree-dump-times {mapper_binding\(struct Y,omp declare mapper ~1Y\) mapper_binding\(struct Z,omp declare mapper ~1Z\)} 2 "original" { target c++ } } } */
|
||||
/* { dg-final { scan-tree-dump {mapper_binding\(struct Z,omp declare mapper ~1Z\) mapper_binding\(struct P,omp declare mapper ~1P\)} "original" { target c++ } } } */
|
||||
|
||||
/* { dg-final { scan-tree-dump {mapper_binding\(struct Z,#pragma omp declare mapper \(struct Z z\) map\(tofrom:z\)\) mapper_binding\(struct Y,#pragma omp declare mapper \(struct Y y\) map\(tofrom:y\)\)} "original" { target c } } } */
|
||||
/* { dg-final { scan-tree-dump {mapper_binding\(struct Z,#pragma omp declare mapper \(struct Z z\) map\(tofrom:z\.z\)\) mapper_binding\(struct Y,#pragma omp declare mapper \(struct Y y\) map\(tofrom:y\.z\)\)} "original" { target c } } } */
|
||||
/* { dg-final { scan-tree-dump {mapper_binding\(struct P,#pragma omp declare mapper \(struct P x\) map\(tofrom:\(x\.zp\)\[0:1\]\) map\(to:x.zp\)\) mapper_binding\(struct Z,#pragma omp declare mapper \(struct Z z\) map\(tofrom:z\.z\)\)} "original" { target c } } } */
|
||||
|
|
@ -0,0 +1,26 @@
|
|||
/* { dg-do compile { target c++ } } */
|
||||
|
||||
typedef struct S_ {
|
||||
int *myarr;
|
||||
int size;
|
||||
} S;
|
||||
|
||||
#pragma omp declare mapper (named: struct S_ v) map(to:v.size, v.myarr) \
|
||||
map(tofrom: v.myarr[0:v.size])
|
||||
/* { dg-error "previous '#pragma omp declare mapper'" "" { target c } .-2 } */
|
||||
/* { dg-note "'#pragma omp declare mapper \\(named: S_\\)' previously defined here" "" { target c++ } .-3 } */
|
||||
|
||||
#pragma omp declare mapper (named: S v) map(to:v.size, v.myarr) \
|
||||
map(tofrom: v.myarr[0:v.size])
|
||||
/* { dg-error "redeclaration of 'named' '#pragma omp declare mapper' for type 'S' \\\{aka 'struct S_'\\\}" "" { target c } .-2 } */
|
||||
/* { dg-error "redefinition of '#pragma omp declare mapper \\(named: S\\)'" "" { target c++ } .-3 } */
|
||||
|
||||
#pragma omp declare mapper (struct S_ v) map(to:v.size, v.myarr) \
|
||||
map(tofrom: v.myarr[0:v.size])
|
||||
/* { dg-error "previous '#pragma omp declare mapper'" "" { target c } .-2 } */
|
||||
/* { dg-note "'#pragma omp declare mapper \\(S_\\)' previously defined here" "" { target c++ } .-3 } */
|
||||
|
||||
#pragma omp declare mapper (S v) map(to:v.size, v.myarr) \
|
||||
map(tofrom: v.myarr[0:v.size])
|
||||
/* { dg-error "redeclaration of '<default>' '#pragma omp declare mapper' for type 'S' \\\{aka 'struct S_'\\\}" "" { target c } .-2 } */
|
||||
/* { dg-error "redefinition of '#pragma omp declare mapper \\(S\\)'" "" { target c++ } .-3 } */
|
||||
|
|
@ -0,0 +1,23 @@
|
|||
/* { dg-do compile { target c++ } } */
|
||||
|
||||
int x = 5;
|
||||
|
||||
struct Q {
|
||||
int *arr1;
|
||||
int *arr2;
|
||||
int *arr3;
|
||||
};
|
||||
|
||||
#pragma omp declare mapper (struct Q myq) map(myq.arr2[0:x])
|
||||
|
||||
struct R {
|
||||
int *arr1;
|
||||
int *arr2;
|
||||
int *arr3;
|
||||
};
|
||||
|
||||
#pragma omp declare mapper (struct R myr) map(myr.arr3[0:y])
|
||||
/* { dg-error "'y' undeclared" "" { target c } .-1 } */
|
||||
/* { dg-error "'y' was not declared in this scope" "" { target c++ } .-2 } */
|
||||
|
||||
int y = 7;
|
||||
|
|
@ -0,0 +1,29 @@
|
|||
/* { dg-do compile { target c++ } } */
|
||||
|
||||
struct Q {
|
||||
int *arr1;
|
||||
int *arr2;
|
||||
int *arr3;
|
||||
};
|
||||
|
||||
int foo (void)
|
||||
{
|
||||
int x = 5;
|
||||
#pragma omp declare mapper (struct Q myq) map(myq.arr2[0:x])
|
||||
return x;
|
||||
}
|
||||
|
||||
struct R {
|
||||
int *arr1;
|
||||
int *arr2;
|
||||
int *arr3;
|
||||
};
|
||||
|
||||
int bar (void)
|
||||
{
|
||||
#pragma omp declare mapper (struct R myr) map(myr.arr3[0:y])
|
||||
/* { dg-error "'y' undeclared" "" { target c } .-1 } */
|
||||
/* { dg-error "'y' was not declared in this scope" "" { target c++ } .-2 } */
|
||||
int y = 7;
|
||||
return y;
|
||||
}
|
||||
|
|
@ -0,0 +1,43 @@
|
|||
/* { dg-do compile { target c++ } } */
|
||||
|
||||
struct Q {
|
||||
int *arr1;
|
||||
int *arr2;
|
||||
int *arr3;
|
||||
int len;
|
||||
};
|
||||
|
||||
struct R {
|
||||
struct Q qarr[5];
|
||||
};
|
||||
|
||||
struct R2 {
|
||||
struct Q *qptr;
|
||||
};
|
||||
|
||||
#pragma omp declare mapper (struct Q myq) map(myq.arr1[0:myq.len]) \
|
||||
map(myq.arr2[0:myq.len]) \
|
||||
map(myq.arr3[0:myq.len])
|
||||
|
||||
#pragma omp declare mapper (struct R myr) map(myr.qarr[2:3])
|
||||
|
||||
#pragma omp declare mapper (struct R2 myr2) map(myr2.qptr[2:3])
|
||||
|
||||
int main (int argc, char *argv[])
|
||||
{
|
||||
struct R r;
|
||||
struct R2 r2;
|
||||
int N = 256;
|
||||
|
||||
#pragma omp target
|
||||
/* { dg-message "sorry, unimplemented: user-defined mapper with non-unit length array section" "" { target *-*-* } .-1 } */
|
||||
{
|
||||
for (int i = 2; i < 5; i++)
|
||||
for (int j = 0; j < N; j++)
|
||||
{
|
||||
r.qarr[i].arr1[j]++;
|
||||
r2.qptr[i].arr2[j]++;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
|
@ -0,0 +1,34 @@
|
|||
/* { dg-do compile { target c++ } } */
|
||||
|
||||
int x = 5;
|
||||
|
||||
struct Q {
|
||||
int *arr1;
|
||||
int *arr2;
|
||||
int *arr3;
|
||||
};
|
||||
|
||||
int y = 5;
|
||||
|
||||
#pragma omp declare mapper (struct Q myq) map(myq.arr2[0:x])
|
||||
/* { dg-error "previous '#pragma omp declare mapper'" "" { target c } .-1 } */
|
||||
/* { dg-note "'#pragma omp declare mapper \\(Q\\)' previously defined here" "" { target c++ } .-2 } */
|
||||
|
||||
#pragma omp declare mapper (struct Q myq) map(myq.arr2[0:y])
|
||||
/* { dg-error "redeclaration of '<default>' '#pragma omp declare mapper' for type 'struct Q'" "" { target c } .-1 } */
|
||||
/* { dg-error "redefinition of '#pragma omp declare mapper \\(Q\\)'" "" { target c++ } .-2 } */
|
||||
|
||||
struct R {
|
||||
int *arr1;
|
||||
};
|
||||
|
||||
void foo (void)
|
||||
{
|
||||
#pragma omp declare mapper (struct R myr) map(myr.arr1[0:x])
|
||||
/* { dg-error "previous '#pragma omp declare mapper'" "" { target c } .-1 } */
|
||||
/* { dg-note "'#pragma omp declare mapper \\(R\\)' previously declared here" "" { target c++ } .-2 } */
|
||||
|
||||
#pragma omp declare mapper (struct R myr) map(myr.arr1[0:y])
|
||||
/* { dg-error "redeclaration of '<default>' '#pragma omp declare mapper' for type 'struct R'" "" { target c } .-1 } */
|
||||
/* { dg-error "redeclaration of '#pragma omp declare mapper \\(R\\)'" "" { target c++ } .-2 } */
|
||||
}
|
||||
|
|
@ -13,20 +13,20 @@ foo (void)
|
|||
#pragma omp target map (to:a)
|
||||
;
|
||||
|
||||
#pragma omp target map (a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */
|
||||
;
|
||||
#pragma omp target map (a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'mapper' or 'present'" "" { target c++ } } */
|
||||
; /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present' before 'a'" "" { target c } .-1 } */
|
||||
|
||||
#pragma omp target map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */
|
||||
;
|
||||
#pragma omp target map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'mapper' or 'present'" "" { target c++ } } */
|
||||
; /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present' before 'a'" "" { target c } .-1 } */
|
||||
|
||||
#pragma omp target enter data map(b7) map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */
|
||||
;
|
||||
#pragma omp target enter data map(b7) map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'mapper' or 'present'" "" { target c++ } } */
|
||||
; /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present' before 'a'" "" { target c } .-1 } */
|
||||
|
||||
#pragma omp target exit data map(b7) map (close, a from: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */
|
||||
;
|
||||
#pragma omp target exit data map(b7) map (close, a from: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'mapper' or 'present'" "" { target c++ } } */
|
||||
; /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present' before 'a'" "" { target c } .-1 } */
|
||||
|
||||
#pragma omp target data map(b7) map (close, a from: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */
|
||||
;
|
||||
#pragma omp target data map(b7) map (close, a from: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'mapper' or 'present'" "" { target c++ } } */
|
||||
; /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present' before 'a'" "" { target c } .-1 } */
|
||||
|
||||
|
||||
#pragma omp target map (close a) /* { dg-error "'close' undeclared" "" { target c } } */
|
||||
|
|
|
|||
|
|
@ -0,0 +1,58 @@
|
|||
// { dg-do compile }
|
||||
// { dg-additional-options "-fdump-tree-gimple" }
|
||||
|
||||
// "omp declare mapper" support -- check expansion in gimple.
|
||||
|
||||
struct S {
|
||||
int *ptr;
|
||||
int size;
|
||||
};
|
||||
|
||||
#define N 64
|
||||
|
||||
#pragma omp declare mapper (S w) map(w.size, w.ptr, w.ptr[:w.size])
|
||||
#pragma omp declare mapper (foo:S w) map(to:w.size, w.ptr) map(w.ptr[:w.size])
|
||||
|
||||
int main (int argc, char *argv[])
|
||||
{
|
||||
S s;
|
||||
s.ptr = new int[N];
|
||||
s.size = N;
|
||||
|
||||
#pragma omp declare mapper (bar:S w) map(w.size, w.ptr, w.ptr[:w.size])
|
||||
|
||||
#pragma omp target
|
||||
{
|
||||
for (int i = 0; i < N; i++)
|
||||
s.ptr[i]++;
|
||||
}
|
||||
|
||||
#pragma omp target map(tofrom: s)
|
||||
{
|
||||
for (int i = 0; i < N; i++)
|
||||
s.ptr[i]++;
|
||||
}
|
||||
|
||||
#pragma omp target map(mapper(default), tofrom: s)
|
||||
{
|
||||
for (int i = 0; i < N; i++)
|
||||
s.ptr[i]++;
|
||||
}
|
||||
|
||||
#pragma omp target map(mapper(foo), alloc: s)
|
||||
{
|
||||
for (int i = 0; i < N; i++)
|
||||
s.ptr[i]++;
|
||||
}
|
||||
|
||||
#pragma omp target map(mapper(bar), tofrom: s)
|
||||
{
|
||||
for (int i = 0; i < N; i++)
|
||||
s.ptr[i]++;
|
||||
}
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
// { dg-final { scan-tree-dump-times {map\(struct:s \[len: 2\]\) map\(alloc:s\.ptr \[len: [0-9]+\]\) map\(tofrom:s\.size \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:s\.ptr \[bias: 0\]\)} 4 "gimple" } }
|
||||
// { dg-final { scan-tree-dump-times {map\(struct:s \[len: 2\]\) map\(alloc:s\.ptr \[len: [0-9]+\]\) map\(to:s\.size \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:s\.ptr \[bias: 0\]\)} 1 "gimple" } }
|
||||
|
|
@ -0,0 +1,30 @@
|
|||
// { dg-do compile }
|
||||
|
||||
// Error-checking tests for "omp declare mapper".
|
||||
|
||||
struct S {
|
||||
int *ptr;
|
||||
int size;
|
||||
};
|
||||
|
||||
struct Z {
|
||||
int z;
|
||||
};
|
||||
|
||||
int main (int argc, char *argv[])
|
||||
{
|
||||
#pragma omp declare mapper (S v) map(v.size, v.ptr[:v.size]) // { dg-note "'#pragma omp declare mapper \\(S\\)' previously declared here" }
|
||||
|
||||
/* This one's a duplicate. */
|
||||
#pragma omp declare mapper (default: S v) map (to: v.size) map (v) // { dg-error "redeclaration of '#pragma omp declare mapper \\(S\\)'" }
|
||||
|
||||
/* ...and this one doesn't use a "base language identifier" for the mapper
|
||||
name. */
|
||||
#pragma omp declare mapper (case: S v) map (to: v.size) // { dg-error "expected identifier or 'default' before 'case'" }
|
||||
// { dg-error "expected ':' before 'case'" "" { target *-*-* } .-1 }
|
||||
|
||||
/* A non-struct/class/union type isn't supposed to work. */
|
||||
#pragma omp declare mapper (name:Z [5]foo) map (foo[0].z) // { dg-error "'Z \\\[5\\\]' is not a struct, union or class type in '#pragma omp declare mapper'" }
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
|
@ -0,0 +1,39 @@
|
|||
#pragma omp declare mapper (int v) // { dg-error "missing 'map' clause before end of line" }
|
||||
#pragma omp declare mapper (float v) map() // { dg-error "expected primary-expression before '\\)' token" }
|
||||
// { dg-error "'float' is not a struct, union or class type in '#pragma omp declare mapper'" "" { target *-*-* } .-1 }
|
||||
|
||||
#pragma omp declare mapper (char v) map(v) // { dg-error "'char' is not a struct, union or class type in '#pragma omp declare mapper'" }
|
||||
|
||||
struct XT {
|
||||
int x;
|
||||
};
|
||||
#pragma omp declare mapper (XT y) map() // { dg-error "expected primary-expression before '\\)' token" }
|
||||
|
||||
struct t {
|
||||
int x;
|
||||
};
|
||||
|
||||
typedef struct t myStruct;
|
||||
|
||||
#pragma omp declare mapper(t) // { dg-error "expected unqualified-id before '\\)' token" }
|
||||
#pragma omp declare mapper(struct t) // { dg-error "expected unqualified-id before '\\)' token" }
|
||||
#pragma omp declare mapper(myStruct) // { dg-error "expected unqualified-id before '\\)' token" }
|
||||
|
||||
#pragma omp declare mapper(name : t v) map() // { dg-error "expected primary-expression before '\\)' token" }
|
||||
|
||||
#pragma omp declare mapper(fancy : struct t v) map(always,present,close,mapper(d),tofrom: v) // { dg-error "in 'declare mapper' directives, parameter to 'mapper' modifier must be 'default'" }
|
||||
|
||||
#pragma omp declare mapper(myStruct v) map(v, v.x) // { dg-note "'#pragma omp declare mapper \\(myStruct\\)' previously declared here" }
|
||||
#pragma omp declare mapper(default : t v) map(v, v.x) // { dg-error "redefinition of '#pragma omp declare mapper \\(t\\)'" }
|
||||
|
||||
|
||||
class A { };
|
||||
class B : public virtual A { };
|
||||
|
||||
#pragma omp declare mapper(class B ci) map(ci) // { dg-error "'B' must not be a virtual base class in '#pragma omp declare mapper'" }
|
||||
|
||||
#pragma omp declare mapper(T v) mapper(v) // { dg-error "'T' does not name a type" }
|
||||
|
||||
union u_t { };
|
||||
|
||||
#pragma omp declare mapper(u_t v) map() // { dg-error "expected primary-expression before '\\)' token" }
|
||||
|
|
@ -368,6 +368,10 @@ enum omp_clause_code {
|
|||
/* OpenMP clause: doacross ({source,sink}:vec). */
|
||||
OMP_CLAUSE_DOACROSS,
|
||||
|
||||
/* OpenMP mapper binding: record implicit mappers in scope for aggregate
|
||||
types used within an offload region. */
|
||||
OMP_CLAUSE__MAPPER_BINDING_,
|
||||
|
||||
/* Internal structure to hold OpenACC cache directive's variable-list.
|
||||
#pragma acc cache (variable-list). */
|
||||
OMP_CLAUSE__CACHE_,
|
||||
|
|
|
|||
|
|
@ -1122,6 +1122,15 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
|
|||
case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
|
||||
pp_string (pp, "always,present,tofrom");
|
||||
break;
|
||||
case GOMP_MAP_UNSET:
|
||||
pp_string (pp, "unset");
|
||||
break;
|
||||
case GOMP_MAP_PUSH_MAPPER_NAME:
|
||||
pp_string (pp, "push_mapper");
|
||||
break;
|
||||
case GOMP_MAP_POP_MAPPER_NAME:
|
||||
pp_string (pp, "pop_mapper");
|
||||
break;
|
||||
default:
|
||||
gcc_unreachable ();
|
||||
}
|
||||
|
|
@ -1200,6 +1209,23 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
|
|||
spc, flags, false);
|
||||
goto print_clause_size;
|
||||
|
||||
case OMP_CLAUSE__MAPPER_BINDING_:
|
||||
pp_string (pp, "mapper_binding(");
|
||||
if (OMP_CLAUSE__MAPPER_BINDING__ID (clause))
|
||||
{
|
||||
dump_generic_node (pp, OMP_CLAUSE__MAPPER_BINDING__ID (clause), spc,
|
||||
flags, false);
|
||||
pp_comma (pp);
|
||||
}
|
||||
dump_generic_node (pp,
|
||||
TREE_TYPE (OMP_CLAUSE__MAPPER_BINDING__DECL (clause)),
|
||||
spc, flags, false);
|
||||
pp_comma (pp);
|
||||
dump_generic_node (pp, OMP_CLAUSE__MAPPER_BINDING__MAPPER (clause), spc,
|
||||
flags, false);
|
||||
pp_right_paren (pp);
|
||||
break;
|
||||
|
||||
case OMP_CLAUSE_NUM_TEAMS:
|
||||
pp_string (pp, "num_teams(");
|
||||
if (OMP_CLAUSE_NUM_TEAMS_LOWER_EXPR (clause))
|
||||
|
|
@ -4263,6 +4289,21 @@ dump_generic_node (pretty_printer *pp, tree node, int spc, dump_flags_t flags,
|
|||
pp_string (pp, ")>");
|
||||
break;
|
||||
|
||||
case OMP_DECLARE_MAPPER:
|
||||
pp_string (pp, "#pragma omp declare mapper (");
|
||||
if (OMP_DECLARE_MAPPER_ID (node))
|
||||
{
|
||||
dump_generic_node (pp, OMP_DECLARE_MAPPER_ID (node), spc, flags,
|
||||
false);
|
||||
pp_colon (pp);
|
||||
}
|
||||
dump_generic_node (pp, TREE_TYPE (node), spc, flags, false);
|
||||
pp_space (pp);
|
||||
dump_generic_node (pp, OMP_DECLARE_MAPPER_DECL (node), spc, flags, false);
|
||||
pp_right_paren (pp);
|
||||
dump_omp_clauses (pp, OMP_DECLARE_MAPPER_CLAUSES (node), spc, flags);
|
||||
break;
|
||||
|
||||
case TRANSACTION_EXPR:
|
||||
if (TRANSACTION_EXPR_OUTER (node))
|
||||
pp_string (pp, "__transaction_atomic [[outer]]");
|
||||
|
|
|
|||
|
|
@ -326,6 +326,7 @@ unsigned const char omp_clause_num_ops[] =
|
|||
2, /* OMP_CLAUSE_MAP */
|
||||
1, /* OMP_CLAUSE_HAS_DEVICE_ADDR */
|
||||
1, /* OMP_CLAUSE_DOACROSS */
|
||||
3, /* OMP_CLAUSE__MAPPER_BINDING_ */
|
||||
2, /* OMP_CLAUSE__CACHE_ */
|
||||
1, /* OMP_CLAUSE_DESTROY */
|
||||
2, /* OMP_CLAUSE_INIT */
|
||||
|
|
@ -428,6 +429,7 @@ const char * const omp_clause_code_name[] =
|
|||
"map",
|
||||
"has_device_addr",
|
||||
"doacross",
|
||||
"_mapper_binding_",
|
||||
"_cache_",
|
||||
"destroy",
|
||||
"init",
|
||||
|
|
|
|||
|
|
@ -1346,6 +1346,13 @@ DEFTREECODE (OMP_STRUCTURED_BLOCK, "omp_structured_block", tcc_statement, 1)
|
|||
Operand 0: OMP_MASTER_BODY: Master section body. */
|
||||
DEFTREECODE (OMP_MASTER, "omp_master", tcc_statement, 1)
|
||||
|
||||
/* OpenMP - #pragma omp declare mapper ([id:] type var) [clause1 ... clauseN]
|
||||
Operand 0: Identifier.
|
||||
Operand 1: Variable decl.
|
||||
Operand 2: List of clauses.
|
||||
The type of the construct is used for the type to be mapped. */
|
||||
DEFTREECODE (OMP_DECLARE_MAPPER, "omp_declare_mapper", tcc_statement, 3)
|
||||
|
||||
/* OpenACC - #pragma acc cache (variable1 ... variableN)
|
||||
Operand 0: OACC_CACHE_CLAUSES: List of variables (transformed into
|
||||
OMP_CLAUSE__CACHE_ clauses). */
|
||||
|
|
|
|||
19
gcc/tree.h
19
gcc/tree.h
|
|
@ -1625,6 +1625,13 @@ class auto_suppress_location_wrappers
|
|||
#define OMP_METADIRECTIVE_VARIANT_BODY(v) \
|
||||
TREE_VALUE (TREE_VALUE (v))
|
||||
|
||||
#define OMP_DECLARE_MAPPER_ID(NODE) \
|
||||
TREE_OPERAND (OMP_DECLARE_MAPPER_CHECK (NODE), 0)
|
||||
#define OMP_DECLARE_MAPPER_DECL(NODE) \
|
||||
TREE_OPERAND (OMP_DECLARE_MAPPER_CHECK (NODE), 1)
|
||||
#define OMP_DECLARE_MAPPER_CLAUSES(NODE) \
|
||||
TREE_OPERAND (OMP_DECLARE_MAPPER_CHECK (NODE), 2)
|
||||
|
||||
#define OMP_SCAN_BODY(NODE) TREE_OPERAND (OMP_SCAN_CHECK (NODE), 0)
|
||||
#define OMP_SCAN_CLAUSES(NODE) TREE_OPERAND (OMP_SCAN_CHECK (NODE), 1)
|
||||
|
||||
|
|
@ -2118,6 +2125,18 @@ class auto_suppress_location_wrappers
|
|||
#define OMP_TARGET_DEVICE_MATCHES_PROPERTIES(NODE) \
|
||||
TREE_OPERAND (OMP_TARGET_DEVICE_MATCHES_CHECK (NODE), 1)
|
||||
|
||||
#define OMP_CLAUSE__MAPPER_BINDING__ID(NODE) \
|
||||
OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, \
|
||||
OMP_CLAUSE__MAPPER_BINDING_), 0)
|
||||
|
||||
#define OMP_CLAUSE__MAPPER_BINDING__DECL(NODE) \
|
||||
OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, \
|
||||
OMP_CLAUSE__MAPPER_BINDING_), 1)
|
||||
|
||||
#define OMP_CLAUSE__MAPPER_BINDING__MAPPER(NODE) \
|
||||
OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, \
|
||||
OMP_CLAUSE__MAPPER_BINDING_), 2)
|
||||
|
||||
/* SSA_NAME accessors. */
|
||||
|
||||
/* Whether SSA_NAME NODE is a virtual operand. This simply caches the
|
||||
|
|
|
|||
|
|
@ -209,7 +209,13 @@ enum gomp_map_kind
|
|||
GOMP_MAP_PRESENT_ALLOC = (GOMP_MAP_LAST | 4),
|
||||
GOMP_MAP_PRESENT_TO = (GOMP_MAP_LAST | 5),
|
||||
GOMP_MAP_PRESENT_FROM = (GOMP_MAP_LAST | 6),
|
||||
GOMP_MAP_PRESENT_TOFROM = (GOMP_MAP_LAST | 7)
|
||||
GOMP_MAP_PRESENT_TOFROM = (GOMP_MAP_LAST | 7),
|
||||
/* Unset, used for "declare mapper" maps with no explicit data movement
|
||||
specified. These use the movement specified at the invocation site. */
|
||||
GOMP_MAP_UNSET = (GOMP_MAP_LAST | 8),
|
||||
/* Used to record the name of a named mapper. */
|
||||
GOMP_MAP_PUSH_MAPPER_NAME = (GOMP_MAP_LAST | 9),
|
||||
GOMP_MAP_POP_MAPPER_NAME = (GOMP_MAP_LAST | 10)
|
||||
};
|
||||
|
||||
#define GOMP_MAP_COPY_TO_P(X) \
|
||||
|
|
|
|||
|
|
@ -0,0 +1,87 @@
|
|||
// { dg-do run }
|
||||
|
||||
#include <cstdlib>
|
||||
#include <cassert>
|
||||
|
||||
#define N 64
|
||||
|
||||
struct points
|
||||
{
|
||||
double *x;
|
||||
double *y;
|
||||
double *z;
|
||||
size_t len;
|
||||
};
|
||||
|
||||
#pragma omp declare mapper(points p) map(to:p.x, p.y, p.z) \
|
||||
map(p.x[0:p.len]) \
|
||||
map(p.y[0:p.len]) \
|
||||
map(p.z[0:p.len])
|
||||
|
||||
struct shape
|
||||
{
|
||||
points tmp;
|
||||
points *pts;
|
||||
int metadata[128];
|
||||
};
|
||||
|
||||
#pragma omp declare mapper(shape s) map(tofrom:s.pts, *s.pts) map(alloc:s.tmp)
|
||||
|
||||
void
|
||||
alloc_points (points *pts, size_t sz)
|
||||
{
|
||||
pts->x = new double[sz];
|
||||
pts->y = new double[sz];
|
||||
pts->z = new double[sz];
|
||||
pts->len = sz;
|
||||
for (int i = 0; i < sz; i++)
|
||||
pts->x[i] = pts->y[i] = pts->z[i] = 0;
|
||||
}
|
||||
|
||||
int main (int argc, char *argv[])
|
||||
{
|
||||
shape myshape;
|
||||
points mypts;
|
||||
|
||||
myshape.pts = &mypts;
|
||||
|
||||
alloc_points (&myshape.tmp, N);
|
||||
myshape.pts = new points;
|
||||
alloc_points (myshape.pts, N);
|
||||
|
||||
#pragma omp target map(myshape)
|
||||
{
|
||||
for (int i = 0; i < N; i++)
|
||||
{
|
||||
myshape.pts->x[i]++;
|
||||
myshape.pts->y[i]++;
|
||||
myshape.pts->z[i]++;
|
||||
}
|
||||
}
|
||||
|
||||
for (int i = 0; i < N; i++)
|
||||
{
|
||||
assert (myshape.pts->x[i] == 1);
|
||||
assert (myshape.pts->y[i] == 1);
|
||||
assert (myshape.pts->z[i] == 1);
|
||||
}
|
||||
|
||||
#pragma omp target
|
||||
{
|
||||
for (int i = 0; i < N; i++)
|
||||
{
|
||||
myshape.pts->x[i]++;
|
||||
myshape.pts->y[i]++;
|
||||
myshape.pts->z[i]++;
|
||||
}
|
||||
}
|
||||
|
||||
for (int i = 0; i < N; i++)
|
||||
{
|
||||
assert (myshape.pts->x[i] == 2);
|
||||
assert (myshape.pts->y[i] == 2);
|
||||
assert (myshape.pts->z[i] == 2);
|
||||
}
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
|
@ -0,0 +1,55 @@
|
|||
// { dg-do run }
|
||||
|
||||
#include <cassert>
|
||||
|
||||
#define N 256
|
||||
|
||||
struct doublebuf
|
||||
{
|
||||
int buf_a[N][N];
|
||||
int buf_b[N][N];
|
||||
};
|
||||
|
||||
#pragma omp declare mapper(lo:doublebuf b) map(b.buf_a[0:N/2][0:N]) \
|
||||
map(b.buf_b[0:N/2][0:N])
|
||||
|
||||
#pragma omp declare mapper(hi:doublebuf b) map(b.buf_a[N/2:N/2][0:N]) \
|
||||
map(b.buf_b[N/2:N/2][0:N])
|
||||
|
||||
int main (int argc, char *argv[])
|
||||
{
|
||||
doublebuf db;
|
||||
|
||||
for (int i = 0; i < N; i++)
|
||||
for (int j = 0; j < N; j++)
|
||||
db.buf_a[i][j] = db.buf_b[i][j] = 0;
|
||||
|
||||
#pragma omp target map(mapper(lo), tofrom:db)
|
||||
{
|
||||
for (int i = 0; i < N / 2; i++)
|
||||
for (int j = 0; j < N; j++)
|
||||
{
|
||||
db.buf_a[i][j]++;
|
||||
db.buf_b[i][j]++;
|
||||
}
|
||||
}
|
||||
|
||||
#pragma omp target map(mapper(hi), tofrom:db)
|
||||
{
|
||||
for (int i = N / 2; i < N; i++)
|
||||
for (int j = 0; j < N; j++)
|
||||
{
|
||||
db.buf_a[i][j]++;
|
||||
db.buf_b[i][j]++;
|
||||
}
|
||||
}
|
||||
|
||||
for (int i = 0; i < N; i++)
|
||||
for (int j = 0; j < N; j++)
|
||||
{
|
||||
assert (db.buf_a[i][j] == 1);
|
||||
assert (db.buf_b[i][j] == 1);
|
||||
}
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
|
@ -0,0 +1,63 @@
|
|||
// { dg-do run }
|
||||
|
||||
#include <cstdlib>
|
||||
#include <cassert>
|
||||
|
||||
struct S {
|
||||
int *myarr;
|
||||
};
|
||||
|
||||
#pragma omp declare mapper (S s) map(to:s.myarr) map (tofrom: s.myarr[0:20])
|
||||
|
||||
namespace A {
|
||||
#pragma omp declare mapper (S s) map(to:s.myarr) map (tofrom: s.myarr[0:100])
|
||||
}
|
||||
|
||||
namespace B {
|
||||
#pragma omp declare mapper (S s) map(to:s.myarr) map (tofrom: s.myarr[100:100])
|
||||
}
|
||||
|
||||
namespace A
|
||||
{
|
||||
void incr_a (S my_s)
|
||||
{
|
||||
#pragma omp target
|
||||
{
|
||||
for (int i = 0; i < 100; i++)
|
||||
my_s.myarr[i]++;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
namespace B
|
||||
{
|
||||
void incr_b (S my_s)
|
||||
{
|
||||
#pragma omp target
|
||||
{
|
||||
for (int i = 100; i < 200; i++)
|
||||
my_s.myarr[i]++;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
int main (int argc, char *argv[])
|
||||
{
|
||||
S my_s;
|
||||
|
||||
my_s.myarr = (int *) calloc (200, sizeof (int));
|
||||
|
||||
#pragma omp target
|
||||
{
|
||||
for (int i = 0; i < 20; i++)
|
||||
my_s.myarr[i]++;
|
||||
}
|
||||
|
||||
A::incr_a (my_s);
|
||||
B::incr_b (my_s);
|
||||
|
||||
for (int i = 0; i < 200; i++)
|
||||
assert (my_s.myarr[i] == (i < 20) ? 2 : 1);
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
|
@ -0,0 +1,63 @@
|
|||
// { dg-do run }
|
||||
|
||||
#include <cstdlib>
|
||||
#include <cassert>
|
||||
|
||||
struct S {
|
||||
int *myarr;
|
||||
};
|
||||
|
||||
#pragma omp declare mapper (S s) map(to:s.myarr) map (tofrom: s.myarr[0:20])
|
||||
|
||||
namespace A {
|
||||
#pragma omp declare mapper (S s) map(to:s.myarr) map (tofrom: s.myarr[0:100])
|
||||
}
|
||||
|
||||
namespace B {
|
||||
#pragma omp declare mapper (S s) map(to:s.myarr) map (tofrom: s.myarr[100:100])
|
||||
}
|
||||
|
||||
namespace A
|
||||
{
|
||||
void incr_a (S &my_s)
|
||||
{
|
||||
#pragma omp target
|
||||
{
|
||||
for (int i = 0; i < 100; i++)
|
||||
my_s.myarr[i]++;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
namespace B
|
||||
{
|
||||
void incr_b (S &my_s)
|
||||
{
|
||||
#pragma omp target
|
||||
{
|
||||
for (int i = 100; i < 200; i++)
|
||||
my_s.myarr[i]++;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
int main (int argc, char *argv[])
|
||||
{
|
||||
S my_s;
|
||||
|
||||
my_s.myarr = (int *) calloc (200, sizeof (int));
|
||||
|
||||
#pragma omp target
|
||||
{
|
||||
for (int i = 0; i < 20; i++)
|
||||
my_s.myarr[i]++;
|
||||
}
|
||||
|
||||
A::incr_a (my_s);
|
||||
B::incr_b (my_s);
|
||||
|
||||
for (int i = 0; i < 200; i++)
|
||||
assert (my_s.myarr[i] == (i < 20) ? 2 : 1);
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
|
@ -0,0 +1,52 @@
|
|||
// { dg-do run }
|
||||
|
||||
#include <cassert>
|
||||
|
||||
struct S
|
||||
{
|
||||
int *myarr;
|
||||
int len;
|
||||
};
|
||||
|
||||
class C
|
||||
{
|
||||
S smemb;
|
||||
#pragma omp declare mapper (custom:S s) map(to:s.myarr) \
|
||||
map(tofrom:s.myarr[0:s.len])
|
||||
|
||||
public:
|
||||
C(int l)
|
||||
{
|
||||
smemb.myarr = new int[l];
|
||||
smemb.len = l;
|
||||
for (int i = 0; i < l; i++)
|
||||
smemb.myarr[i] = 0;
|
||||
}
|
||||
void bump();
|
||||
void check();
|
||||
};
|
||||
|
||||
void
|
||||
C::bump ()
|
||||
{
|
||||
#pragma omp target map(mapper(custom), tofrom: smemb)
|
||||
{
|
||||
for (int i = 0; i < smemb.len; i++)
|
||||
smemb.myarr[i]++;
|
||||
}
|
||||
}
|
||||
|
||||
void
|
||||
C::check ()
|
||||
{
|
||||
for (int i = 0; i < smemb.len; i++)
|
||||
assert (smemb.myarr[i] == 1);
|
||||
}
|
||||
|
||||
int main (int argc, char *argv[])
|
||||
{
|
||||
C test (100);
|
||||
test.bump ();
|
||||
test.check ();
|
||||
return 0;
|
||||
}
|
||||
|
|
@ -0,0 +1,37 @@
|
|||
// { dg-do run }
|
||||
|
||||
#include <cassert>
|
||||
|
||||
template <typename T>
|
||||
void adjust (T param)
|
||||
{
|
||||
#pragma omp declare mapper (T x) map(to:x.len, x.base) \
|
||||
map(tofrom:x.base[0:x.len])
|
||||
|
||||
#pragma omp target
|
||||
for (int i = 0; i < param.len; i++)
|
||||
param.base[i]++;
|
||||
}
|
||||
|
||||
struct S {
|
||||
int len;
|
||||
int *base;
|
||||
};
|
||||
|
||||
int main (int argc, char *argv[])
|
||||
{
|
||||
S a;
|
||||
|
||||
a.len = 100;
|
||||
a.base = new int[a.len];
|
||||
|
||||
for (int i = 0; i < a.len; i++)
|
||||
a.base[i] = 0;
|
||||
|
||||
adjust (a);
|
||||
|
||||
for (int i = 0; i < a.len; i++)
|
||||
assert (a.base[i] == 1);
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
|
@ -0,0 +1,59 @@
|
|||
// { dg-do run }
|
||||
|
||||
#include <cassert>
|
||||
|
||||
struct S
|
||||
{
|
||||
int *myarr;
|
||||
};
|
||||
|
||||
struct T
|
||||
{
|
||||
S *s;
|
||||
};
|
||||
|
||||
#pragma omp declare mapper (s100: S x) map(to: x.myarr) \
|
||||
map(tofrom: x.myarr[0:100])
|
||||
// Define this because ...
|
||||
#pragma omp declare mapper (default: S x) map(to: x.myarr) \
|
||||
map(tofrom: x.myarr[0:100])
|
||||
|
||||
|
||||
void
|
||||
bump (T t)
|
||||
{
|
||||
/* Here we have an implicit/default mapper invoking a named mapper. We
|
||||
need to make sure that can be located properly at gimplification
|
||||
time. */
|
||||
|
||||
// ... the following is invalid in OpenMP - albeit supported by GCC
|
||||
// (after disabling: error: in ‘declare mapper’ directives, parameter to ‘mapper’ modifier must be ‘default’ )
|
||||
|
||||
// #pragma omp declare mapper (T t) map(to:t.s) map(mapper(s100), tofrom: t.s[0])
|
||||
|
||||
// ... thus, we now use ...
|
||||
#pragma omp declare mapper (T t) map(to:t.s) map(mapper(default), tofrom: t.s[0])
|
||||
|
||||
#pragma omp target
|
||||
for (int i = 0; i < 100; i++)
|
||||
t.s->myarr[i]++;
|
||||
}
|
||||
|
||||
int main (int argc, char *argv[])
|
||||
{
|
||||
S my_s;
|
||||
T my_t;
|
||||
|
||||
my_s.myarr = new int[100];
|
||||
my_t.s = &my_s;
|
||||
|
||||
for (int i = 0; i < 100; i++)
|
||||
my_s.myarr[i] = 0;
|
||||
|
||||
bump (my_t);
|
||||
|
||||
for (int i = 0; i < 100; i++)
|
||||
assert (my_s.myarr[i] == 1);
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
|
@ -0,0 +1,61 @@
|
|||
// { dg-do run }
|
||||
|
||||
#include <cassert>
|
||||
|
||||
struct S
|
||||
{
|
||||
int *myarr;
|
||||
int len;
|
||||
};
|
||||
|
||||
template<typename T>
|
||||
class C
|
||||
{
|
||||
T memb;
|
||||
#pragma omp declare mapper (T t) map(to:t.len, t.myarr) \
|
||||
map(tofrom:t.myarr[0:t.len])
|
||||
|
||||
public:
|
||||
C(int sz);
|
||||
~C();
|
||||
void bump();
|
||||
void check();
|
||||
};
|
||||
|
||||
template<typename T>
|
||||
C<T>::C(int sz)
|
||||
{
|
||||
memb.myarr = new int[sz];
|
||||
for (int i = 0; i < sz; i++)
|
||||
memb.myarr[i] = 0;
|
||||
memb.len = sz;
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
C<T>::~C()
|
||||
{
|
||||
delete[] memb.myarr;
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void C<T>::bump()
|
||||
{
|
||||
#pragma omp target map(memb)
|
||||
for (int i = 0; i < memb.len; i++)
|
||||
memb.myarr[i]++;
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void C<T>::check()
|
||||
{
|
||||
for (int i = 0; i < memb.len; i++)
|
||||
assert (memb.myarr[i] == 1);
|
||||
}
|
||||
|
||||
int main(int argc, char *argv[])
|
||||
{
|
||||
C<S> c_int(100);
|
||||
c_int.bump();
|
||||
c_int.check();
|
||||
return 0;
|
||||
}
|
||||
|
|
@ -0,0 +1,64 @@
|
|||
/* { dg-do run { target c++ } } */
|
||||
|
||||
#include <string.h>
|
||||
#include <stdlib.h>
|
||||
#include <assert.h>
|
||||
|
||||
#define N 64
|
||||
|
||||
typedef struct {
|
||||
int *arr;
|
||||
int size;
|
||||
} B;
|
||||
|
||||
#pragma omp declare mapper (mapB : B myb) map(to: myb.size, myb.arr) \
|
||||
map(tofrom: myb.arr[0:myb.size])
|
||||
// While GCC handles more, only default is ...
|
||||
#pragma omp declare mapper (default : B myb) map(to: myb.size, myb.arr) \
|
||||
map(tofrom: myb.arr[0:myb.size])
|
||||
|
||||
struct A {
|
||||
int *arr1;
|
||||
B *arr2;
|
||||
int arr3[N];
|
||||
};
|
||||
|
||||
int
|
||||
main (int argc, char *argv[])
|
||||
{
|
||||
struct A var;
|
||||
|
||||
memset (&var, 0, sizeof var);
|
||||
var.arr1 = (int *) calloc (N, sizeof (int));
|
||||
var.arr2 = (B *) malloc (sizeof (B));
|
||||
var.arr2->arr = (int *) calloc (N, sizeof (float));
|
||||
var.arr2->size = N;
|
||||
|
||||
{
|
||||
// ... permitted here:
|
||||
#pragma omp declare mapper (struct A x) map(to: x.arr1, x.arr2) \
|
||||
map(tofrom: x.arr1[0:N]) \
|
||||
map(mapper(default), tofrom: x.arr2[0:1])
|
||||
#pragma omp target
|
||||
{
|
||||
for (int i = 0; i < N; i++)
|
||||
{
|
||||
var.arr1[i]++;
|
||||
var.arr2->arr[i]++;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
for (int i = 0; i < N; i++)
|
||||
{
|
||||
assert (var.arr1[i] == 1);
|
||||
assert (var.arr2->arr[i] == 1);
|
||||
assert (var.arr3[i] == 0);
|
||||
}
|
||||
|
||||
free (var.arr1);
|
||||
free (var.arr2->arr);
|
||||
free (var.arr2);
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
|
@ -0,0 +1,59 @@
|
|||
/* { dg-do run { target c++ } } */
|
||||
|
||||
#include <string.h>
|
||||
#include <stdlib.h>
|
||||
#include <assert.h>
|
||||
|
||||
#define N 64
|
||||
|
||||
typedef struct B_tag {
|
||||
int *arr;
|
||||
int size;
|
||||
} B;
|
||||
|
||||
#pragma omp declare mapper (B myb) map(to: myb.size, myb.arr) \
|
||||
map(tofrom: myb.arr[0:myb.size])
|
||||
|
||||
struct A {
|
||||
int *arr1;
|
||||
B *arr2;
|
||||
int arr3[N];
|
||||
};
|
||||
|
||||
int
|
||||
main (int argc, char *argv[])
|
||||
{
|
||||
struct A var;
|
||||
|
||||
memset (&var, 0, sizeof var);
|
||||
var.arr1 = (int *) calloc (N, sizeof (int));
|
||||
var.arr2 = (B *) malloc (sizeof (B));
|
||||
var.arr2->arr = (int *) calloc (N, sizeof (int));
|
||||
var.arr2->size = N;
|
||||
|
||||
{
|
||||
#pragma omp declare mapper (struct A x) map(to: x.arr1, x.arr2) \
|
||||
map(tofrom: x.arr1[0:N]) map(tofrom: x.arr2[0:1])
|
||||
#pragma omp target
|
||||
{
|
||||
for (int i = 0; i < N; i++)
|
||||
{
|
||||
var.arr1[i]++;
|
||||
var.arr2->arr[i]++;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
for (int i = 0; i < N; i++)
|
||||
{
|
||||
assert (var.arr1[i] == 1);
|
||||
assert (var.arr2->arr[i] == 1);
|
||||
assert (var.arr3[i] == 0);
|
||||
}
|
||||
|
||||
free (var.arr1);
|
||||
free (var.arr2->arr);
|
||||
free (var.arr2);
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
|
@ -0,0 +1,94 @@
|
|||
/* { dg-do run { target c++ } } */
|
||||
|
||||
#include <string.h>
|
||||
#include <stdlib.h>
|
||||
#include <assert.h>
|
||||
|
||||
#define N 64
|
||||
|
||||
typedef struct {
|
||||
int *arr;
|
||||
int size;
|
||||
} B;
|
||||
|
||||
#pragma omp declare mapper (samename : B myb) map(to: myb.size, myb.arr) \
|
||||
map(tofrom: myb.arr[0:myb.size])
|
||||
// While GCC handles more, only default is ...
|
||||
#pragma omp declare mapper (default : B myb) map(to: myb.size, myb.arr) \
|
||||
map(tofrom: myb.arr[0:myb.size])
|
||||
typedef struct {
|
||||
int *arr;
|
||||
int size;
|
||||
} C;
|
||||
|
||||
|
||||
struct A {
|
||||
int *arr1;
|
||||
B *arr2;
|
||||
C *arr3;
|
||||
};
|
||||
|
||||
int
|
||||
main (int argc, char *argv[])
|
||||
{
|
||||
struct A var;
|
||||
|
||||
memset (&var, 0, sizeof var);
|
||||
var.arr1 = (int *) calloc (N, sizeof (int));
|
||||
var.arr2 = (B *) malloc (sizeof (B));
|
||||
var.arr2->arr = (int *) calloc (N, sizeof (int));
|
||||
var.arr2->size = N;
|
||||
var.arr3 = (C *) malloc (sizeof (C));
|
||||
var.arr3->arr = (int *) calloc (N, sizeof (int));
|
||||
var.arr3->size = N;
|
||||
|
||||
{
|
||||
// ... permitted here.
|
||||
#pragma omp declare mapper (struct A x) map(to: x.arr1, x.arr2) \
|
||||
map(tofrom: x.arr1[0:N]) \
|
||||
map(mapper(default), tofrom: x.arr2[0:1])
|
||||
#pragma omp target
|
||||
{
|
||||
for (int i = 0; i < N; i++)
|
||||
{
|
||||
var.arr1[i]++;
|
||||
var.arr2->arr[i]++;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
{
|
||||
#pragma omp declare mapper (samename : C myc) map(to: myc.size, myc.arr) \
|
||||
map(tofrom: myc.arr[0:myc.size])
|
||||
// While GCC handles more, only default is ...
|
||||
#pragma omp declare mapper (default : C myc) map(to: myc.size, myc.arr) \
|
||||
map(tofrom: myc.arr[0:myc.size])
|
||||
// ... permitted here.
|
||||
#pragma omp declare mapper (struct A x) map(to: x.arr1, x.arr3) \
|
||||
map(tofrom: x.arr1[0:N]) \
|
||||
map(mapper( default ) , tofrom: *x.arr3)
|
||||
#pragma omp target
|
||||
{
|
||||
for (int i = 0; i < N; i++)
|
||||
{
|
||||
var.arr1[i]++;
|
||||
var.arr3->arr[i]++;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
for (int i = 0; i < N; i++)
|
||||
{
|
||||
assert (var.arr1[i] == 2);
|
||||
assert (var.arr2->arr[i] == 1);
|
||||
assert (var.arr3->arr[i] == 1);
|
||||
}
|
||||
|
||||
free (var.arr1);
|
||||
free (var.arr2->arr);
|
||||
free (var.arr2);
|
||||
free (var.arr3->arr);
|
||||
free (var.arr3);
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
|
@ -0,0 +1,55 @@
|
|||
/* { dg-do run { target c++ } } */
|
||||
|
||||
#include <assert.h>
|
||||
|
||||
struct T {
|
||||
int a;
|
||||
int b;
|
||||
int c;
|
||||
};
|
||||
|
||||
void foo (void)
|
||||
{
|
||||
struct T x;
|
||||
x.a = x.b = x.c = 0;
|
||||
|
||||
#pragma omp target
|
||||
{
|
||||
x.a++;
|
||||
x.c++;
|
||||
}
|
||||
|
||||
assert (x.a == 1);
|
||||
assert (x.b == 0);
|
||||
assert (x.c == 1);
|
||||
}
|
||||
|
||||
// An identity mapper. This should do the same thing as the default!
|
||||
#pragma omp declare mapper (struct T v) map(v)
|
||||
|
||||
void bar (void)
|
||||
{
|
||||
struct T x;
|
||||
x.a = x.b = x.c = 0;
|
||||
|
||||
#pragma omp target
|
||||
{
|
||||
x.b++;
|
||||
}
|
||||
|
||||
#pragma omp target map(x)
|
||||
{
|
||||
x.a++;
|
||||
}
|
||||
|
||||
assert (x.a == 1);
|
||||
assert (x.b == 1);
|
||||
assert (x.c == 0);
|
||||
}
|
||||
|
||||
int main (int argc, char *argv[])
|
||||
{
|
||||
foo ();
|
||||
bar ();
|
||||
return 0;
|
||||
}
|
||||
|
|
@ -0,0 +1,57 @@
|
|||
/* { dg-do run { target c++ } } */
|
||||
|
||||
#include <stdlib.h>
|
||||
#include <assert.h>
|
||||
|
||||
struct Z {
|
||||
int *arr;
|
||||
};
|
||||
|
||||
void baz (struct Z *zarr, int len)
|
||||
{
|
||||
#pragma omp declare mapper (struct Z myvar) map(to: myvar.arr) \
|
||||
map(tofrom: myvar.arr[0:len])
|
||||
zarr[0].arr = (int *) calloc (len, sizeof (int));
|
||||
zarr[5].arr = (int *) calloc (len, sizeof (int));
|
||||
|
||||
#pragma omp target map(zarr, *zarr)
|
||||
{
|
||||
for (int i = 0; i < len; i++)
|
||||
zarr[0].arr[i]++;
|
||||
}
|
||||
|
||||
#pragma omp target map(zarr, zarr[5])
|
||||
{
|
||||
for (int i = 0; i < len; i++)
|
||||
zarr[5].arr[i]++;
|
||||
}
|
||||
|
||||
#pragma omp target map(zarr[5])
|
||||
{
|
||||
for (int i = 0; i < len; i++)
|
||||
zarr[5].arr[i]++;
|
||||
}
|
||||
|
||||
#pragma omp target map(zarr, zarr[5:1])
|
||||
{
|
||||
for (int i = 0; i < len; i++)
|
||||
zarr[5].arr[i]++;
|
||||
}
|
||||
|
||||
for (int i = 0; i < len; i++)
|
||||
assert (zarr[0].arr[i] == 1);
|
||||
|
||||
for (int i = 0; i < len; i++)
|
||||
assert (zarr[5].arr[i] == 3);
|
||||
|
||||
free (zarr[5].arr);
|
||||
free (zarr[0].arr);
|
||||
}
|
||||
|
||||
int
|
||||
main (int argc, char *argv[])
|
||||
{
|
||||
struct Z myzarr[10];
|
||||
baz (myzarr, 256);
|
||||
return 0;
|
||||
}
|
||||
|
|
@ -0,0 +1,62 @@
|
|||
/* { dg-do run { target c++ } } */
|
||||
|
||||
#include <string.h>
|
||||
#include <stdlib.h>
|
||||
#include <assert.h>
|
||||
|
||||
#define N 64
|
||||
|
||||
struct A {
|
||||
int *arr1;
|
||||
float *arr2;
|
||||
int arr3[N];
|
||||
};
|
||||
|
||||
int
|
||||
main (int argc, char *argv[])
|
||||
{
|
||||
struct A var;
|
||||
|
||||
memset (&var, 0, sizeof var);
|
||||
var.arr1 = (int *) calloc (N, sizeof (int));
|
||||
var.arr2 = (float *) calloc (N, sizeof (float));
|
||||
|
||||
{
|
||||
#pragma omp declare mapper (struct A x) map(to: x.arr1) \
|
||||
map(tofrom: x.arr1[0:N])
|
||||
#pragma omp target
|
||||
{
|
||||
for (int i = 0; i < N; i++)
|
||||
var.arr1[i]++;
|
||||
}
|
||||
}
|
||||
|
||||
{
|
||||
#pragma omp declare mapper (struct A x) map(to: x.arr2) \
|
||||
map(tofrom: x.arr2[0:N])
|
||||
#pragma omp target
|
||||
{
|
||||
for (int i = 0; i < N; i++)
|
||||
var.arr2[i]++;
|
||||
}
|
||||
}
|
||||
|
||||
{
|
||||
#pragma omp declare mapper (struct A x) map(tofrom: x.arr3[0:N])
|
||||
#pragma omp target
|
||||
{
|
||||
for (int i = 0; i < N; i++)
|
||||
var.arr3[i]++;
|
||||
}
|
||||
}
|
||||
|
||||
for (int i = 0; i < N; i++)
|
||||
{
|
||||
assert (var.arr1[i] == 1);
|
||||
assert (var.arr2[i] == 1);
|
||||
assert (var.arr3[i] == 1);
|
||||
}
|
||||
|
||||
free (var.arr1);
|
||||
free (var.arr2);
|
||||
}
|
||||
Loading…
Reference in New Issue