git
/
gcc
mirror of git://gcc.gnu.org/git/gcc.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

OpenMP/OpenACC: Rework clause expansion and nested struct handling 

This patch reworks clause expansion in the C, C++ and (to a lesser
extent) Fortran front ends for OpenMP and OpenACC mapping nodes used in
GPU offloading support.

At present a single clause may be turned into several mapping nodes,
or have its mapping type changed, in several places scattered through
the front- and middle-end.  The analysis relating to which particular
transformations are needed for some given expression has become quite hard
to follow.  Briefly, we manipulate clause types in the following places:

 1. During parsing, in c_omp_adjust_map_clauses.  Depending on a set of
    rules, we may change a FIRSTPRIVATE_POINTER (etc.) mapping into
    ATTACH_DETACH, or mark the decl addressable.

 2. In semantics.cc or c-typeck.cc, clauses are expanded in
    handle_omp_array_sections (called via {c_}finish_omp_clauses, or in
    finish_omp_clauses itself.  The two cases are for processing array
    sections (the former), or non-array sections (the latter).

 3. In gimplify.cc, we build sibling lists for struct accesses, which
    groups and sorts accesses along with their struct base, creating
    new ALLOC/RELEASE nodes for pointers.

 4. In gimplify.cc:gimplify_adjust_omp_clauses, mapping nodes may be
    adjusted or created.

This patch doesn't completely disrupt this scheme, though clause
types are no longer adjusted in c_omp_adjust_map_clauses (step 1).
Clause expansion in step 2 (for C and C++) now uses a single, unified
mechanism, parts of which are also reused for analysis in step 3.

Rather than the kind-of "ad-hoc" pattern matching on addresses used to
expand clauses used at present, a new method for analysing addresses is
introduced.  This does a recursive-descent tree walk on expression nodes,
and emits a vector of tokens describing each "part" of the address.
This tokenized address can then be translated directly into mapping nodes,
with the assurance that no part of the expression has been inadvertently
skipped or misinterpreted.  In this way, all the variations of ways
pointers, arrays, references and component accesses might be combined
can be teased apart into easily-understood cases - and we know we've
"parsed" the whole address before we start analysis, so the right code
paths can easily be selected.

For example, a simple access "arr[idx]" might parse as:

  base-decl access-indexed-array

or "mystruct->foo[x]" with a pointer "foo" component might parse as:

  base-decl access-pointer component-selector access-pointer

A key observation is that support for "array" bases, e.g. accesses
whose root nodes are not structures, but describe scalars or arrays,
and also *one-level deep* structure accesses, have first-class support
in gimplify and beyond.  Expressions that use deeper struct accesses
or e.g. multiple indirections were more problematic: some cases worked,
but lots of cases didn't.  This patch reimplements the support for those
in gimplify.cc, again using the new "address tokenization" support.

An expression like "mystruct->foo->bar[0:10]" used in a mapping node will
translate the right-hand access directly in the front-end.  The base for
the access will be "mystruct->foo".  This is handled recursively in
gimplify.cc -- there may be several accesses of "mystruct"'s members
on the same directive, so the sibling-list building machinery can be
used again.  (This was already being done for OpenACC, but the new
implementation differs somewhat in details, and is more robust.)

For OpenMP, in the case where the base pointer itself,
i.e. "mystruct->foo" here, is NOT mapped on the same directive, we
create a "fragile" mapping.  This turns the "foo" component access
into a zero-length allocation (which is a new feature for the runtime,
so support has been added there too).

A couple of changes have been made to how mapping clauses are turned
into mapping nodes:

The first change is based on the observation that it is probably never
correct to use GOMP_MAP_ALWAYS_POINTER for component accesses (e.g. for
references), because if the containing struct is already mapped on the
target then the host version of the pointer in question will be corrupted
if the struct is copied back from the target.  This patch removes all
such uses, across each of C, C++ and Fortran.

The second change is to the way that GOMP_MAP_ATTACH_DETACH nodes
are processed during sibling-list creation.  For OpenMP, for pointer
components, we must map the base pointer separately from an array section
that uses the base pointer, so e.g. we must have both "map(mystruct.base)"
and "map(mystruct.base[0:10])" mappings.  These create nodes such as:

  GOMP_MAP_TOFROM mystruct.base
  G_M_TOFROM *mystruct.base [len: 10*elemsize] G_M_ATTACH_DETACH mystruct.base

Instead of using the first of these directly when building the struct
sibling list then skipping the group using GOMP_MAP_ATTACH_DETACH,
leading to:

  GOMP_MAP_STRUCT mystruct [len: 1] GOMP_MAP_TOFROM mystruct.base

we now introduce a new "mini-pass", omp_resolve_clause_dependencies, that
drops the GOMP_MAP_TOFROM for the base pointer, marks the second group
as having had a base-pointer mapping, then omp_build_struct_sibling_lists
can create:

  GOMP_MAP_STRUCT mystruct [len: 1] GOMP_MAP_ALLOC mystruct.base [len: ptrsize]

This ends up working better in many cases, particularly those involving
references.  (The "alloc" space is immediately overwritten by a pointer
attachment, so this is mildly more efficient than a redundant TO mapping
at runtime also.)

There is support in the address tokenizer for "arbitrary" base expressions
which aren't rooted at a decl, but that is not used as present because
such addresses are disallowed at parse time.

In the front-ends, the address tokenization machinery is mostly only
used for clause expansion and not for diagnostics at present.  It could
be used for those too, which would allow more of my previous "address
inspector" implementation to be removed.

The new bits in gimplify.cc work with OpenACC also.

This version of the patch addresses several first-pass review comments
from Tobias, and fixes a few previously-missed cases for manually-managed
ragged array mappings (including cases using references).  Some arbitrary
differences between handling of clause expansion for C vs. C++ have also
been fixed, and some fragments from later in the patch series have been
moved forward (where they were useful for fixing bugs).  Several new
test cases have been added.

2023-11-29  Julian Brown  <julian@codesourcery.com>

gcc/c-family/
	* c-common.h (c_omp_region_type): Add C_ORT_EXIT_DATA,
	C_ORT_OMP_EXIT_DATA and C_ORT_ACC_TARGET.
	(omp_addr_token): Add forward declaration.
	(c_omp_address_inspector): New class.
	* c-omp.cc (c_omp_adjust_map_clauses): Mark decls addressable here, but
	do not change any mapping node types.
	(c_omp_address_inspector::unconverted_ref_origin,
	c_omp_address_inspector::component_access_p,
	c_omp_address_inspector::check_clause,
	c_omp_address_inspector::get_root_term,
	c_omp_address_inspector::map_supported_p,
	c_omp_address_inspector::get_origin,
	c_omp_address_inspector::maybe_unconvert_ref,
	c_omp_address_inspector::maybe_zero_length_array_section,
	c_omp_address_inspector::expand_array_base,
	c_omp_address_inspector::expand_component_selector,
	c_omp_address_inspector::expand_map_clause): New methods.
	(omp_expand_access_chain): New function.

gcc/c/
	* c-parser.cc (c_parser_oacc_all_clauses): Add TARGET_P parameter. Use
	to select region type for c_finish_omp_clauses call.
	(c_parser_oacc_loop): Update calls to c_parser_oacc_all_clauses.
	(c_parser_oacc_compute): Likewise.
	(c_parser_omp_target_data, c_parser_omp_target_enter_data): Support
	ATTACH kind.
	(c_parser_omp_target_exit_data): Support DETACH kind.
	(check_clauses): Handle GOMP_MAP_POINTER and GOMP_MAP_ATTACH here.
	* c-typeck.cc (handle_omp_array_sections_1,
	handle_omp_array_sections, c_finish_omp_clauses): Use
	c_omp_address_inspector class and OMP address tokenizer to analyze and
	expand map clause expressions.  Fix some diagnostics.  Fix "is OpenACC"
	condition for C_ORT_ACC_TARGET addition.

gcc/cp/
	* parser.cc (cp_parser_oacc_all_clauses): Add TARGET_P parameter. Use
	to select region type for finish_omp_clauses call.
	(cp_parser_omp_target_data, cp_parser_omp_target_enter_data): Support
	GOMP_MAP_ATTACH kind.
	(cp_parser_omp_target_exit_data): Support GOMP_MAP_DETACH kind.
	(cp_parser_oacc_declare): Update call to cp_parser_oacc_all_clauses.
	(cp_parser_oacc_loop): Update calls to cp_parser_oacc_all_clauses.
	(cp_parser_oacc_compute): Likewise.
	* pt.cc (tsubst_expr): Use C_ORT_ACC_TARGET for call to
	tsubst_omp_clauses for OpenACC compute regions.
	* semantics.cc (cp_omp_address_inspector): New class, derived from
	c_omp_address_inspector.
	(handle_omp_array_sections_1, handle_omp_array_sections,
	finish_omp_clauses): Use cp_omp_address_inspector class and OMP address
	tokenizer to analyze and expand OpenMP map clause expressions.  Fix
	some diagnostics.  Support C_ORT_ACC_TARGET.
	(finish_omp_target): Handle GOMP_MAP_POINTER.

gcc/fortran/
	* trans-openmp.cc (gfc_trans_omp_array_section): Add OPENMP parameter.
	Use GOMP_MAP_ATTACH_DETACH instead of GOMP_MAP_ALWAYS_POINTER for
	derived type components.
	(gfc_trans_omp_clauses): Update calls to gfc_trans_omp_array_section.

gcc/
	* gimplify.cc (build_struct_comp_nodes): Don't process
	GOMP_MAP_ATTACH_DETACH "middle" nodes here.
	(omp_mapping_group): Add REPROCESS_STRUCT and FRAGILE booleans for
	nested struct handling.
	(omp_strip_components_and_deref, omp_strip_indirections): Remove
	functions.
	(omp_get_attachment): Handle GOMP_MAP_DETACH here.
	(omp_group_last): Handle GOMP_MAP_*, GOMP_MAP_DETACH,
	GOMP_MAP_ATTACH_DETACH groups for "exit data" of reference-to-pointer
	component array sections.
	(omp_gather_mapping_groups_1): Initialise reprocess_struct and fragile
	fields.
	(omp_group_base): Handle GOMP_MAP_ATTACH_DETACH after GOMP_MAP_STRUCT.
	(omp_index_mapping_groups_1): Skip reprocess_struct groups.
	(omp_get_nonfirstprivate_group, omp_directive_maps_explicitly,
	omp_resolve_clause_dependencies, omp_first_chained_access_token): New
	functions.
	(omp_check_mapping_compatibility): Adjust accepted node combinations
	for "from" clauses using release instead of alloc.
	(omp_accumulate_sibling_list): Add GROUP_MAP, ADDR_TOKENS, FRAGILE_P,
	REPROCESSING_STRUCT, ADDED_TAIL parameters.  Use OMP address tokenizer
	to analyze addresses.  Reimplement nested struct handling, and
	implement "fragile groups".
	(omp_build_struct_sibling_lists): Adjust for changes to
	omp_accumulate_sibling_list.  Recalculate bias for ATTACH_DETACH nodes
	after GOMP_MAP_STRUCT nodes.
	(gimplify_scan_omp_clauses): Call omp_resolve_clause_dependencies.  Use
	OMP address tokenizer.
	(gimplify_adjust_omp_clauses_1): Use build_fold_indirect_ref_loc
	instead of build_simple_mem_ref_loc.
	* omp-general.cc (omp-general.h, tree-pretty-print.h): Include.
	(omp_addr_tokenizer): New namespace.
	(omp_addr_tokenizer::omp_addr_token): New.
	(omp_addr_tokenizer::omp_parse_component_selector,
	omp_addr_tokenizer::omp_parse_ref,
	omp_addr_tokenizer::omp_parse_pointer,
	omp_addr_tokenizer::omp_parse_access_method,
	omp_addr_tokenizer::omp_parse_access_methods,
	omp_addr_tokenizer::omp_parse_structure_base,
	omp_addr_tokenizer::omp_parse_structured_expr,
	omp_addr_tokenizer::omp_parse_array_expr,
	omp_addr_tokenizer::omp_access_chain_p,
	omp_addr_tokenizer::omp_accessed_addr): New functions.
	(omp_parse_expr, debug_omp_tokenized_addr): New functions.
	* omp-general.h (omp_addr_tokenizer::access_method_kinds,
	omp_addr_tokenizer::structure_base_kinds,
	omp_addr_tokenizer::token_type,
	omp_addr_tokenizer::omp_addr_token,
	omp_addr_tokenizer::omp_access_chain_p,
	omp_addr_tokenizer::omp_accessed_addr): New.
	(omp_addr_token, omp_parse_expr): New.
	* omp-low.cc (scan_sharing_clauses): Skip error check for references
	to pointers.
	* tree.h (OMP_CLAUSE_ATTACHMENT_MAPPING_ERASED): New macro.

gcc/testsuite/
	* c-c++-common/gomp/clauses-2.c: Fix error output.
	* c-c++-common/gomp/target-implicit-map-2.c: Adjust scan output.
	* c-c++-common/gomp/target-50.c: Adjust scan output.
	* c-c++-common/gomp/target-enter-data-1.c: Adjust scan output.
	* g++.dg/gomp/static-component-1.C: New test.
	* gcc.dg/gomp/target-3.c: Adjust scan output.
	* gfortran.dg/gomp/map-9.f90: Adjust scan output.

libgomp/
	* target.c (gomp_map_pointer): Modify zero-length array section
	pointer handling.
	(gomp_attach_pointer): Likewise.
	(gomp_map_fields_existing): Use gomp_map_0len_lookup.
	(gomp_attach_pointer): Allow attaching null pointers (or Fortran
	"unassociated" pointers).
	(gomp_map_vars_internal): Handle zero-sized struct members.  Add
	diagnostic for unmapped struct pointer members.
	* testsuite/libgomp.c-c++-common/baseptrs-1.c: New test.
	* testsuite/libgomp.c-c++-common/baseptrs-2.c: New test.
	* testsuite/libgomp.c-c++-common/baseptrs-6.c: New test.
	* testsuite/libgomp.c-c++-common/baseptrs-7.c: New test.
	* testsuite/libgomp.c-c++-common/ptr-attach-2.c: New test.
	* testsuite/libgomp.c-c++-common/target-implicit-map-2.c: Fix missing
	"free".
	* testsuite/libgomp.c-c++-common/target-implicit-map-5.c: New test.
	* testsuite/libgomp.c-c++-common/target-map-zlas-1.c: New test.
	* testsuite/libgomp.c++/class-array-1.C: New test.
	* testsuite/libgomp.c++/baseptrs-3.C: New test.
	* testsuite/libgomp.c++/baseptrs-4.C: New test.
	* testsuite/libgomp.c++/baseptrs-5.C: New test.
	* testsuite/libgomp.c++/baseptrs-8.C: New test.
	* testsuite/libgomp.c++/baseptrs-9.C: New test.
	* testsuite/libgomp.c++/ref-mapping-1.C: New test.
	* testsuite/libgomp.c++/target-48.C: New test.
	* testsuite/libgomp.c++/target-49.C: New test.
	* testsuite/libgomp.c++/target-exit-data-reftoptr-1.C: New test.
	* testsuite/libgomp.c++/target-lambda-1.C: Update for OpenMP 5.2
	semantics.
	* testsuite/libgomp.c++/target-this-3.C: Likewise.
	* testsuite/libgomp.c++/target-this-4.C: Likewise.
	* testsuite/libgomp.fortran/struct-elem-map-1.f90: Add temporary XFAIL.
	* testsuite/libgomp.fortran/target-enter-data-6.f90: Likewise.
This commit is contained in:
Julian Brown 2023-08-14 12:41:56 +00:00
parent e1fde9de3f
commit 5fdb150cd4
44 changed files with 7216 additions and 820 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

71
gcc/c-family/c-common.h
View File

@ -1279,8 +1279,11 @@ enum c_omp_region_type
C_ORT_ACC = 1 << 1,
C_ORT_DECLARE_SIMD = 1 << 2,
C_ORT_TARGET = 1 << 3,
C_ORT_EXIT_DATA = 1 << 4,
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_TARGET = C_ORT_OMP | C_ORT_TARGET,
C_ORT_OMP_EXIT_DATA = C_ORT_OMP | C_ORT_EXIT_DATA,
C_ORT_ACC_TARGET = C_ORT_ACC | C_ORT_TARGET
};
extern tree c_finish_omp_master (location_t, tree);
@ -1317,6 +1320,72 @@ 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);
namespace omp_addr_tokenizer { struct omp_addr_token; }
typedef omp_addr_tokenizer::omp_addr_token omp_addr_token;
class c_omp_address_inspector
{
location_t loc;
tree root_term;
bool indirections;
int map_supported;
protected:
tree orig;
public:
c_omp_address_inspector (location_t loc, tree t)
: loc (loc), root_term (NULL_TREE), indirections (false),
map_supported (-1), orig (t)
{
}
~c_omp_address_inspector ()
{
}
virtual bool processing_template_decl_p ()
{
return false;
}
virtual void emit_unmappable_type_notes (tree)
{
}
virtual tree convert_from_reference (tree)
{
gcc_unreachable ();
}
virtual tree build_array_ref (location_t loc, tree arr, tree idx)
{
tree eltype = TREE_TYPE (TREE_TYPE (arr));
return build4_loc (loc, ARRAY_REF, eltype, arr, idx, NULL_TREE,
NULL_TREE);
}
virtual bool check_clause (tree);
tree get_root_term (bool);
tree unconverted_ref_origin ();
bool component_access_p ();
bool map_supported_p ();
static tree get_origin (tree);
static tree maybe_unconvert_ref (tree);
bool maybe_zero_length_array_section (tree);
tree expand_array_base (tree, vec<omp_addr_token *> &, tree, unsigned *,
c_omp_region_type);
tree expand_component_selector (tree, vec<omp_addr_token *> &, tree,
unsigned *, c_omp_region_type);
tree expand_map_clause (tree, tree, vec<omp_addr_token *> &,
c_omp_region_type);
};
enum c_omp_directive_kind {
C_OMP_DIR_STANDALONE,
C_OMP_DIR_CONSTRUCT,

875
gcc/c-family/c-omp.cc
View File

@ -3169,8 +3169,9 @@ struct map_clause
decl_mapped (false), omp_declare_target (false) { }
};
/* Adjust map clauses after normal clause parsing, mainly to turn specific
base-pointer map cases into attach/detach and mark them addressable. */
/* Adjust map clauses after normal clause parsing, mainly to mark specific
base-pointer map cases addressable that may be turned into attach/detach
operations during gimplification. */
void
c_omp_adjust_map_clauses (tree clauses, bool is_target)
{
@ -3186,7 +3187,6 @@ c_omp_adjust_map_clauses (tree clauses, bool is_target)
&& POINTER_TYPE_P (TREE_TYPE (OMP_CLAUSE_DECL (c))))
{
tree ptr = OMP_CLAUSE_DECL (c);
OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ATTACH_DETACH);
c_common_mark_addressable_vec (ptr);
}
return;
@ -3199,7 +3199,7 @@ c_omp_adjust_map_clauses (tree clauses, bool is_target)
&& DECL_P (OMP_CLAUSE_DECL (c)))
{
/* If this is for a target construct, the firstprivate pointer
is changed to attach/detach if either is true:
is marked addressable if either is true:
(1) the base-pointer is mapped in this same construct, or
(2) the base-pointer is a variable place on the device by
"declare target" directives.
@ -3241,11 +3241,874 @@ c_omp_adjust_map_clauses (tree clauses, bool is_target)
if (mc.firstprivate_ptr_p
&& (mc.decl_mapped || mc.omp_declare_target))
c_common_mark_addressable_vec (OMP_CLAUSE_DECL (mc.clause));
}
}
/* Maybe strip off an indirection from a "converted" reference, then find the
origin of a pointer (i.e. without any offset). */
tree
c_omp_address_inspector::unconverted_ref_origin ()
{
tree t = orig;
/* We may have a reference-typed component access at the outermost level
that has had convert_from_reference called on it. Get the un-dereferenced
reference itself. */
t = maybe_unconvert_ref (t);
/* Find base pointer for POINTER_PLUS_EXPR, etc. */
t = get_origin (t);
return t;
}
/* Return TRUE if the address is a component access. */
bool
c_omp_address_inspector::component_access_p ()
{
tree t = maybe_unconvert_ref (orig);
t = get_origin (t);
return TREE_CODE (t) == COMPONENT_REF;
}
/* Perform various checks on the address, as described by clause CLAUSE (we
only use its code and location here). */
bool
c_omp_address_inspector::check_clause (tree clause)
{
tree t = unconverted_ref_origin ();
if (TREE_CODE (t) != COMPONENT_REF)
return true;
if (TREE_CODE (TREE_OPERAND (t, 1)) == FIELD_DECL
&& DECL_BIT_FIELD (TREE_OPERAND (t, 1)))
{
error_at (OMP_CLAUSE_LOCATION (clause),
"bit-field %qE in %qs clause",
t, omp_clause_code_name[OMP_CLAUSE_CODE (clause)]);
return false;
}
else if (!processing_template_decl_p ()
&& !omp_mappable_type (TREE_TYPE (t)))
{
error_at (OMP_CLAUSE_LOCATION (clause),
"%qE does not have a mappable type in %qs clause",
t, omp_clause_code_name[OMP_CLAUSE_CODE (clause)]);
emit_unmappable_type_notes (TREE_TYPE (t));
return false;
}
else if (TREE_TYPE (t) && TYPE_ATOMIC (TREE_TYPE (t)))
{
error_at (OMP_CLAUSE_LOCATION (clause),
"%<_Atomic%> %qE in %qs clause", t,
omp_clause_code_name[OMP_CLAUSE_CODE (clause)]);
return false;
}
return true;
}
/* Find the "root term" for the address. This is the innermost decl, etc.
of the access. */
tree
c_omp_address_inspector::get_root_term (bool checking)
{
if (root_term && !checking)
return root_term;
tree t = unconverted_ref_origin ();
while (TREE_CODE (t) == COMPONENT_REF)
{
if (checking
&& TREE_TYPE (TREE_OPERAND (t, 0))
&& TREE_CODE (TREE_TYPE (TREE_OPERAND (t, 0))) == UNION_TYPE)
{
OMP_CLAUSE_SET_MAP_KIND (mc.clause, GOMP_MAP_ATTACH_DETACH);
c_common_mark_addressable_vec (OMP_CLAUSE_DECL (mc.clause));
error_at (loc, "%qE is a member of a union", t);
return error_mark_node;
}
t = TREE_OPERAND (t, 0);
while (TREE_CODE (t) == MEM_REF
|| TREE_CODE (t) == INDIRECT_REF
|| TREE_CODE (t) == ARRAY_REF)
{
if (TREE_CODE (t) == MEM_REF
|| TREE_CODE (t) == INDIRECT_REF)
indirections = true;
t = TREE_OPERAND (t, 0);
STRIP_NOPS (t);
if (TREE_CODE (t) == POINTER_PLUS_EXPR)
t = TREE_OPERAND (t, 0);
}
}
root_term = t;
return t;
}
/* Return TRUE if the address is supported in mapping clauses. At present,
this means that the innermost expression is a DECL_P, but could be extended
to other types of expression in the future. */
bool
c_omp_address_inspector::map_supported_p ()
{
/* If we've already decided if the mapped address is supported, return
that. */
if (map_supported != -1)
return map_supported;
tree t = unconverted_ref_origin ();
STRIP_NOPS (t);
while (TREE_CODE (t) == INDIRECT_REF
|| TREE_CODE (t) == MEM_REF
|| TREE_CODE (t) == ARRAY_REF
|| TREE_CODE (t) == COMPONENT_REF
|| TREE_CODE (t) == COMPOUND_EXPR
|| TREE_CODE (t) == SAVE_EXPR
|| TREE_CODE (t) == POINTER_PLUS_EXPR
|| TREE_CODE (t) == NON_LVALUE_EXPR
|| TREE_CODE (t) == NOP_EXPR)
if (TREE_CODE (t) == COMPOUND_EXPR)
t = TREE_OPERAND (t, 1);
else
t = TREE_OPERAND (t, 0);
STRIP_NOPS (t);
map_supported = DECL_P (t);
return map_supported;
}
/* Get the origin of an address T, stripping off offsets and some other
bits. */
tree
c_omp_address_inspector::get_origin (tree t)
{
while (1)
{
if (TREE_CODE (t) == COMPOUND_EXPR)
{
t = TREE_OPERAND (t, 1);
STRIP_NOPS (t);
}
else if (TREE_CODE (t) == POINTER_PLUS_EXPR
|| TREE_CODE (t) == SAVE_EXPR)
t = TREE_OPERAND (t, 0);
else if (TREE_CODE (t) == INDIRECT_REF
&& TREE_CODE (TREE_TYPE (TREE_OPERAND (t, 0))) == REFERENCE_TYPE)
t = TREE_OPERAND (t, 0);
else
break;
}
STRIP_NOPS (t);
return t;
}
/* For an address T that might be a reference that has had
"convert_from_reference" called on it, return the actual reference without
any indirection. */
tree
c_omp_address_inspector::maybe_unconvert_ref (tree t)
{
if (TREE_CODE (t) == INDIRECT_REF
&& TREE_CODE (TREE_TYPE (TREE_OPERAND (t, 0))) == REFERENCE_TYPE)
return TREE_OPERAND (t, 0);
return t;
}
/* Return TRUE if CLAUSE might describe a zero-length array section. */
bool
c_omp_address_inspector::maybe_zero_length_array_section (tree clause)
{
switch (OMP_CLAUSE_MAP_KIND (clause))
{
case GOMP_MAP_ALLOC:
case GOMP_MAP_IF_PRESENT:
case GOMP_MAP_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_TOFROM:
case GOMP_MAP_ALWAYS_TO:
case GOMP_MAP_ALWAYS_FROM:
case GOMP_MAP_ALWAYS_TOFROM:
case GOMP_MAP_PRESENT_ALLOC:
case GOMP_MAP_PRESENT_TO:
case GOMP_MAP_PRESENT_FROM:
case GOMP_MAP_PRESENT_TOFROM:
case GOMP_MAP_ALWAYS_PRESENT_TO:
case GOMP_MAP_ALWAYS_PRESENT_FROM:
case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
case GOMP_MAP_RELEASE:
case GOMP_MAP_DELETE:
case GOMP_MAP_FORCE_TO:
case GOMP_MAP_FORCE_FROM:
case GOMP_MAP_FORCE_TOFROM:
case GOMP_MAP_FORCE_PRESENT:
return true;
default:
return false;
}
}
/* Expand a chained access. We only expect to see a quite limited range of
expression types here, because e.g. you can't have an array of
references. */
static tree
omp_expand_access_chain (tree c, tree expr, vec<omp_addr_token *> &addr_tokens,
unsigned *idx, c_omp_region_type ort)
{
using namespace omp_addr_tokenizer;
location_t loc = OMP_CLAUSE_LOCATION (c);
unsigned i = *idx;
tree c2 = NULL_TREE;
gomp_map_kind kind;
if ((ort & C_ORT_EXIT_DATA) != 0
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM
|| (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FROM
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DELETE
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_RELEASE
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_FROM
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_FROM
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PRESENT_FROM
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_PRESENT_FROM)))
kind = GOMP_MAP_DETACH;
else
kind = GOMP_MAP_ATTACH;
switch (addr_tokens[i]->u.access_kind)
{
case ACCESS_POINTER:
case ACCESS_POINTER_OFFSET:
{
tree virtual_origin
= fold_convert_loc (loc, ptrdiff_type_node, addr_tokens[i]->expr);
tree data_addr = omp_accessed_addr (addr_tokens, i, expr);
c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (c2, kind);
OMP_CLAUSE_DECL (c2) = addr_tokens[i]->expr;
OMP_CLAUSE_SIZE (c2)
= fold_build2_loc (loc, MINUS_EXPR, ptrdiff_type_node,
fold_convert_loc (loc, ptrdiff_type_node,
data_addr),
virtual_origin);
}
break;
case ACCESS_INDEXED_ARRAY:
break;
default:
return error_mark_node;
}
if (c2)
{
OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
OMP_CLAUSE_CHAIN (c) = c2;
c = c2;
}
*idx = ++i;
if (i < addr_tokens.length ()
&& addr_tokens[i]->type == ACCESS_METHOD)
return omp_expand_access_chain (c, expr, addr_tokens, idx, ort);
return c;
}
/* Translate "array_base_decl access_method" to OMP mapping clauses. */
tree
c_omp_address_inspector::expand_array_base (tree c,
vec<omp_addr_token *> &addr_tokens,
tree expr, unsigned *idx,
c_omp_region_type ort)
{
using namespace omp_addr_tokenizer;
location_t loc = OMP_CLAUSE_LOCATION (c);
int i = *idx;
tree decl = addr_tokens[i + 1]->expr;
bool decl_p = DECL_P (decl);
bool declare_target_p = (decl_p
&& is_global_var (decl)
&& lookup_attribute ("omp declare target",
DECL_ATTRIBUTES (decl)));
bool map_p = OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP;
bool implicit_p = map_p && OMP_CLAUSE_MAP_IMPLICIT (c);
bool chain_p = omp_access_chain_p (addr_tokens, i + 1);
tree c2 = NULL_TREE, c3 = NULL_TREE;
unsigned consume_tokens = 2;
bool target_p = (ort & C_ORT_TARGET) != 0;
bool openmp_p = (ort & C_ORT_OMP) != 0;
gcc_assert (i == 0);
if (!openmp_p
&& map_p
&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
{
i += 2;
*idx = i;
return c;
}
switch (addr_tokens[i + 1]->u.access_kind)
{
case ACCESS_DIRECT:
if (decl_p && !target_p)
c_common_mark_addressable_vec (addr_tokens[i + 1]->expr);
break;
case ACCESS_REF:
{
/* Copy the referenced object. Note that we do this even for !MAP_P
clauses. */
tree obj = convert_from_reference (addr_tokens[i + 1]->expr);
if (TREE_CODE (TREE_TYPE (obj)) == ARRAY_TYPE)
/* We have a ref to array: add a [0] element as the ME expects. */
OMP_CLAUSE_DECL (c) = build_array_ref (loc, obj, integer_zero_node);
else
OMP_CLAUSE_DECL (c) = obj;
OMP_CLAUSE_SIZE (c) = TYPE_SIZE_UNIT (TREE_TYPE (obj));
if (!map_p)
{
if (decl_p)
c_common_mark_addressable_vec (addr_tokens[i + 1]->expr);
break;
}
if (!target_p)
break;
/* If we have a reference to a pointer, avoid using
FIRSTPRIVATE_REFERENCE here in case the pointer is modified in the
offload region (we can only do that if the pointer does not point
to a mapped block). We could avoid doing this if we don't have a
FROM mapping... */
bool ref_to_ptr = TREE_CODE (TREE_TYPE (obj)) == POINTER_TYPE;
c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
if (!ref_to_ptr
&& !declare_target_p
&& decl_p)
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_REFERENCE);
else
{
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
if (decl_p)
c_common_mark_addressable_vec (addr_tokens[i + 1]->expr);
}
OMP_CLAUSE_DECL (c2) = addr_tokens[i + 1]->expr;
OMP_CLAUSE_SIZE (c2) = size_zero_node;
if (ref_to_ptr)
{
c3 = c2;
c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALLOC);
OMP_CLAUSE_DECL (c2) = addr_tokens[i + 1]->expr;
OMP_CLAUSE_SIZE (c2)
= TYPE_SIZE_UNIT (TREE_TYPE (OMP_CLAUSE_DECL (c2)));
}
}
break;
case ACCESS_INDEXED_REF_TO_ARRAY:
{
if (!map_p)
{
if (decl_p)
c_common_mark_addressable_vec (addr_tokens[i + 1]->expr);
break;
}
if (!target_p)
break;
tree virtual_origin
= convert_from_reference (addr_tokens[i + 1]->expr);
virtual_origin = build_fold_addr_expr (virtual_origin);
virtual_origin = fold_convert_loc (loc, ptrdiff_type_node,
virtual_origin);
tree data_addr = omp_accessed_addr (addr_tokens, i + 1, expr);
c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
if (decl_p && target_p && !declare_target_p)
{
/* It appears that omp-low.cc mishandles cases where we have a
[reference to an] array of pointers such as:
int *arr[N]; (or "int *(&arr)[N] = ...")
#pragma omp target map(arr[a][b:c])
{ ... }
in such cases chain_p will be true. For now, fall back to
GOMP_MAP_POINTER. */
enum gomp_map_kind k = chain_p ? GOMP_MAP_POINTER
: GOMP_MAP_FIRSTPRIVATE_REFERENCE;
OMP_CLAUSE_SET_MAP_KIND (c2, k);
}
else
{
if (decl_p)
c_common_mark_addressable_vec (addr_tokens[i + 1]->expr);
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
}
OMP_CLAUSE_DECL (c2) = addr_tokens[i + 1]->expr;
OMP_CLAUSE_SIZE (c2)
= fold_build2_loc (loc, MINUS_EXPR, ptrdiff_type_node,
fold_convert_loc (loc, ptrdiff_type_node,
data_addr),
virtual_origin);
}
break;
case ACCESS_INDEXED_ARRAY:
{
if (!map_p)
{
if (decl_p)
c_common_mark_addressable_vec (addr_tokens[i + 1]->expr);
break;
}
/* The code handling "firstprivatize_array_bases" in gimplify.cc is
relevant here. What do we need to create for arrays at this
stage? (This condition doesn't feel quite right. FIXME?) */
if (!target_p
&& (TREE_CODE (TREE_TYPE (addr_tokens[i + 1]->expr))
== ARRAY_TYPE))
break;
tree virtual_origin
= build_fold_addr_expr (addr_tokens[i + 1]->expr);
virtual_origin = fold_convert_loc (loc, ptrdiff_type_node,
virtual_origin);
tree data_addr = omp_accessed_addr (addr_tokens, i + 1, expr);
c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
if (decl_p && target_p)
{
/* See comment for ACCESS_INDEXED_REF_TO_ARRAY above. */
enum gomp_map_kind k = chain_p ? GOMP_MAP_POINTER
: GOMP_MAP_FIRSTPRIVATE_POINTER;
OMP_CLAUSE_SET_MAP_KIND (c2, k);
}
else
{
if (decl_p)
c_common_mark_addressable_vec (addr_tokens[i + 1]->expr);
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
}
OMP_CLAUSE_DECL (c2) = addr_tokens[i + 1]->expr;
OMP_CLAUSE_SIZE (c2)
= fold_build2_loc (loc, MINUS_EXPR, ptrdiff_type_node,
fold_convert_loc (loc, ptrdiff_type_node,
data_addr),
virtual_origin);
}
break;
case ACCESS_POINTER:
case ACCESS_POINTER_OFFSET:
{
if (!map_p)
{
if (decl_p)
c_common_mark_addressable_vec (addr_tokens[i + 1]->expr);
break;
}
unsigned last_access = i + 1;
tree virtual_origin;
if (chain_p
&& addr_tokens[i + 2]->type == ACCESS_METHOD
&& addr_tokens[i + 2]->u.access_kind == ACCESS_INDEXED_ARRAY)
{
/* !!! This seems wrong for ACCESS_POINTER_OFFSET. */
consume_tokens = 3;
chain_p = omp_access_chain_p (addr_tokens, i + 2);
last_access = i + 2;
virtual_origin
= build_array_ref (loc, addr_tokens[last_access]->expr,
integer_zero_node);
virtual_origin = build_fold_addr_expr (virtual_origin);
virtual_origin = fold_convert_loc (loc, ptrdiff_type_node,
virtual_origin);
}
else
virtual_origin = fold_convert_loc (loc, ptrdiff_type_node,
addr_tokens[last_access]->expr);
tree data_addr = omp_accessed_addr (addr_tokens, last_access, expr);
c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
/* For OpenACC, use FIRSTPRIVATE_POINTER for decls even on non-compute
regions (e.g. "acc data" constructs). It'll be removed anyway in
gimplify.cc, but doing it this way maintains diagnostic
behaviour. */
if (decl_p && (target_p || !openmp_p) && !chain_p && !declare_target_p)
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
else
{
if (decl_p)
c_common_mark_addressable_vec (addr_tokens[i + 1]->expr);
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
}
OMP_CLAUSE_DECL (c2) = addr_tokens[i + 1]->expr;
OMP_CLAUSE_SIZE (c2)
= fold_build2_loc (loc, MINUS_EXPR, ptrdiff_type_node,
fold_convert_loc (loc, ptrdiff_type_node,
data_addr),
virtual_origin);
}
break;
case ACCESS_REF_TO_POINTER:
case ACCESS_REF_TO_POINTER_OFFSET:
{
if (!map_p)
{
if (decl_p)
c_common_mark_addressable_vec (addr_tokens[i + 1]->expr);
break;
}
unsigned last_access = i + 1;
tree virtual_origin;
if (chain_p
&& addr_tokens[i + 2]->type == ACCESS_METHOD
&& addr_tokens[i + 2]->u.access_kind == ACCESS_INDEXED_ARRAY)
{
/* !!! This seems wrong for ACCESS_POINTER_OFFSET. */
consume_tokens = 3;
chain_p = omp_access_chain_p (addr_tokens, i + 2);
last_access = i + 2;
virtual_origin
= build_array_ref (loc, addr_tokens[last_access]->expr,
integer_zero_node);
virtual_origin = build_fold_addr_expr (virtual_origin);
virtual_origin = fold_convert_loc (loc, ptrdiff_type_node,
virtual_origin);
}
else
{
virtual_origin
= convert_from_reference (addr_tokens[last_access]->expr);
virtual_origin = fold_convert_loc (loc, ptrdiff_type_node,
virtual_origin);
}
tree data_addr = omp_accessed_addr (addr_tokens, last_access, expr);
c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
if (decl_p && target_p && !chain_p && !declare_target_p)
{
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_REFERENCE);
OMP_CLAUSE_DECL (c2) = addr_tokens[i + 1]->expr;
}
else
{
if (decl_p)
c_common_mark_addressable_vec (addr_tokens[i + 1]->expr);
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
OMP_CLAUSE_DECL (c2)
= convert_from_reference (addr_tokens[i + 1]->expr);
}
OMP_CLAUSE_SIZE (c2)
= fold_build2_loc (loc, MINUS_EXPR, ptrdiff_type_node,
fold_convert_loc (loc, ptrdiff_type_node,
data_addr),
virtual_origin);
}
break;
default:
*idx = i + consume_tokens;
return error_mark_node;
}
if (c3)
{
OMP_CLAUSE_CHAIN (c3) = OMP_CLAUSE_CHAIN (c);
OMP_CLAUSE_CHAIN (c2) = c3;
OMP_CLAUSE_CHAIN (c) = c2;
if (implicit_p)
{
OMP_CLAUSE_MAP_IMPLICIT (c2) = 1;
OMP_CLAUSE_MAP_IMPLICIT (c3) = 1;
}
c = c3;
}
else if (c2)
{
OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
OMP_CLAUSE_CHAIN (c) = c2;
if (implicit_p)
OMP_CLAUSE_MAP_IMPLICIT (c2) = 1;
c = c2;
}
i += consume_tokens;
*idx = i;
if (chain_p && map_p)
return omp_expand_access_chain (c, expr, addr_tokens, idx, ort);
return c;
}
/* Translate "component_selector access_method" to OMP mapping clauses. */
tree
c_omp_address_inspector::expand_component_selector (tree c,
vec<omp_addr_token *>
&addr_tokens,
tree expr, unsigned *idx,
c_omp_region_type ort)
{
using namespace omp_addr_tokenizer;
location_t loc = OMP_CLAUSE_LOCATION (c);
unsigned i = *idx;
tree c2 = NULL_TREE, c3 = NULL_TREE;
bool chain_p = omp_access_chain_p (addr_tokens, i + 1);
bool map_p = OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP;
switch (addr_tokens[i + 1]->u.access_kind)
{
case ACCESS_DIRECT:
case ACCESS_INDEXED_ARRAY:
break;
case ACCESS_REF:
{
/* Copy the referenced object. Note that we also do this for !MAP_P
clauses. */
tree obj = convert_from_reference (addr_tokens[i + 1]->expr);
OMP_CLAUSE_DECL (c) = obj;
OMP_CLAUSE_SIZE (c) = TYPE_SIZE_UNIT (TREE_TYPE (obj));
if (!map_p)
break;
c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
OMP_CLAUSE_DECL (c2) = addr_tokens[i + 1]->expr;
OMP_CLAUSE_SIZE (c2) = size_zero_node;
}
break;
case ACCESS_INDEXED_REF_TO_ARRAY:
{
if (!map_p)
break;
tree virtual_origin
= convert_from_reference (addr_tokens[i + 1]->expr);
virtual_origin = build_fold_addr_expr (virtual_origin);
virtual_origin = fold_convert_loc (loc, ptrdiff_type_node,
virtual_origin);
tree data_addr = omp_accessed_addr (addr_tokens, i + 1, expr);
c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
OMP_CLAUSE_DECL (c2) = addr_tokens[i + 1]->expr;
OMP_CLAUSE_SIZE (c2)
= fold_build2_loc (loc, MINUS_EXPR, ptrdiff_type_node,
fold_convert_loc (loc, ptrdiff_type_node,
data_addr),
virtual_origin);
}
break;
case ACCESS_POINTER:
case ACCESS_POINTER_OFFSET:
{
if (!map_p)
break;
tree virtual_origin
= fold_convert_loc (loc, ptrdiff_type_node,
addr_tokens[i + 1]->expr);
tree data_addr = omp_accessed_addr (addr_tokens, i + 1, expr);
c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
OMP_CLAUSE_DECL (c2) = addr_tokens[i + 1]->expr;
OMP_CLAUSE_SIZE (c2)
= fold_build2_loc (loc, MINUS_EXPR, ptrdiff_type_node,
fold_convert_loc (loc, ptrdiff_type_node,
data_addr),
virtual_origin);
}
break;
case ACCESS_REF_TO_POINTER:
case ACCESS_REF_TO_POINTER_OFFSET:
{
if (!map_p)
break;
tree ptr = convert_from_reference (addr_tokens[i + 1]->expr);
tree virtual_origin = fold_convert_loc (loc, ptrdiff_type_node,
ptr);
tree data_addr = omp_accessed_addr (addr_tokens, i + 1, expr);
/* Attach the pointer... */
c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
OMP_CLAUSE_DECL (c2) = ptr;
OMP_CLAUSE_SIZE (c2)
= fold_build2_loc (loc, MINUS_EXPR, ptrdiff_type_node,
fold_convert_loc (loc, ptrdiff_type_node,
data_addr),
virtual_origin);
/* ...and also the reference. */
c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (c3, GOMP_MAP_ATTACH_DETACH);
OMP_CLAUSE_DECL (c3) = addr_tokens[i + 1]->expr;
OMP_CLAUSE_SIZE (c3) = size_zero_node;
}
break;
default:
*idx = i + 2;
return error_mark_node;
}
if (c3)
{
OMP_CLAUSE_CHAIN (c3) = OMP_CLAUSE_CHAIN (c);
OMP_CLAUSE_CHAIN (c2) = c3;
OMP_CLAUSE_CHAIN (c) = c2;
c = c3;
}
else if (c2)
{
OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
OMP_CLAUSE_CHAIN (c) = c2;
c = c2;
}
i += 2;
*idx = i;
if (chain_p && map_p)
return omp_expand_access_chain (c, expr, addr_tokens, idx, ort);
return c;
}
/* Expand a map clause into a group of mapping clauses, creating nodes to
attach/detach pointers and so forth as necessary. */
tree
c_omp_address_inspector::expand_map_clause (tree c, tree expr,
vec<omp_addr_token *> &addr_tokens,
c_omp_region_type ort)
{
using namespace omp_addr_tokenizer;
unsigned i, length = addr_tokens.length ();
for (i = 0; i < length;)
{
int remaining = length - i;
if (remaining >= 2
&& addr_tokens[i]->type == ARRAY_BASE
&& addr_tokens[i]->u.structure_base_kind == BASE_DECL
&& addr_tokens[i + 1]->type == ACCESS_METHOD)
{
c = expand_array_base (c, addr_tokens, expr, &i, ort);
if (c == error_mark_node)
return error_mark_node;
}
else if (remaining >= 2
&& addr_tokens[i]->type == ARRAY_BASE
&& addr_tokens[i]->u.structure_base_kind == BASE_ARBITRARY_EXPR
&& addr_tokens[i + 1]->type == ACCESS_METHOD)
{
c = expand_array_base (c, addr_tokens, expr, &i, ort);
if (c == error_mark_node)
return error_mark_node;
}
else if (remaining >= 2
&& addr_tokens[i]->type == STRUCTURE_BASE
&& addr_tokens[i]->u.structure_base_kind == BASE_DECL
&& addr_tokens[i + 1]->type == ACCESS_METHOD)
{
if (addr_tokens[i + 1]->u.access_kind == ACCESS_DIRECT)
c_common_mark_addressable_vec (addr_tokens[i + 1]->expr);
i += 2;
while (addr_tokens[i]->type == ACCESS_METHOD)
i++;
}
else if (remaining >= 2
&& addr_tokens[i]->type == STRUCTURE_BASE
&& addr_tokens[i]->u.structure_base_kind == BASE_ARBITRARY_EXPR
&& addr_tokens[i + 1]->type == ACCESS_METHOD)
{
switch (addr_tokens[i + 1]->u.access_kind)
{
case ACCESS_DIRECT:
case ACCESS_POINTER:
i += 2;
while (addr_tokens[i]->type == ACCESS_METHOD)
i++;
break;
default:
return error_mark_node;
}
}
else if (remaining >= 2
&& addr_tokens[i]->type == COMPONENT_SELECTOR
&& addr_tokens[i + 1]->type == ACCESS_METHOD)
{
c = expand_component_selector (c, addr_tokens, expr, &i, ort);
/* We used 'expr', so these must have been the last tokens. */
gcc_assert (i == length);
if (c == error_mark_node)
return error_mark_node;
}
else if (remaining >= 3
&& addr_tokens[i]->type == COMPONENT_SELECTOR
&& addr_tokens[i + 1]->type == STRUCTURE_BASE
&& (addr_tokens[i + 1]->u.structure_base_kind
== BASE_COMPONENT_EXPR)
&& addr_tokens[i + 2]->type == ACCESS_METHOD)
{
i += 3;
while (addr_tokens[i]->type == ACCESS_METHOD)
i++;
}
else
break;
}
if (i == length)
return c;
return error_mark_node;
}
const struct c_omp_directive c_omp_directives[] = {

23
gcc/c/c-parser.cc
View File

@ -19063,7 +19063,8 @@ c_parser_omp_clause_detach (c_parser *parser, tree list)
static tree
c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
const char *where, bool finish_p = true)
const char *where, bool finish_p = true,
bool target_p = false)
{
tree clauses = NULL;
bool first = true;
@ -19273,7 +19274,8 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
c_parser_skip_to_pragma_eol (parser);
if (finish_p)
return c_finish_omp_clauses (clauses, C_ORT_ACC);
return c_finish_omp_clauses (clauses, target_p ? C_ORT_ACC_TARGET
: C_ORT_ACC);
return clauses;
}
@ -20011,12 +20013,13 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
mask |= OACC_LOOP_CLAUSE_MASK;
tree clauses = c_parser_oacc_all_clauses (parser, mask, p_name,
cclauses == NULL);
/*finish_p=*/cclauses == NULL,
/*target=*/is_parallel);
if (cclauses)
{
clauses = c_oacc_split_loop_clauses (clauses, cclauses, is_parallel);
if (*cclauses)
*cclauses = c_finish_omp_clauses (*cclauses, C_ORT_ACC);
*cclauses = c_finish_omp_clauses (*cclauses, C_ORT_ACC_TARGET);
if (clauses)
clauses = c_finish_omp_clauses (clauses, C_ORT_ACC);
}
@ -20144,7 +20147,9 @@ c_parser_oacc_compute (location_t loc, c_parser *parser,
}
}
tree clauses = c_parser_oacc_all_clauses (parser, mask, p_name);
tree clauses = c_parser_oacc_all_clauses (parser, mask, p_name,
/*finish_p=*/true,
/*target=*/true);
tree block = c_begin_omp_parallel ();
add_stmt (c_parser_omp_structured_block (parser, if_p));
@ -23670,6 +23675,7 @@ c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p)
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_ALWAYS_POINTER:
case GOMP_MAP_ATTACH_DETACH:
case GOMP_MAP_ATTACH:
break;
default:
map_seen |= 1;
@ -23835,6 +23841,7 @@ c_parser_omp_target_enter_data (location_t loc, c_parser *parser,
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_ALWAYS_POINTER:
case GOMP_MAP_ATTACH_DETACH:
case GOMP_MAP_ATTACH:
break;
default:
map_seen |= 1;
@ -23909,7 +23916,8 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser,
tree clauses
= c_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK,
"#pragma omp target exit data");
"#pragma omp target exit data", false);
clauses = c_finish_omp_clauses (clauses, C_ORT_OMP_EXIT_DATA);
c_omp_adjust_map_clauses (clauses, false);
int map_seen = 0;
for (tree *pc = &clauses; *pc;)
@ -23944,6 +23952,7 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser,
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_ALWAYS_POINTER:
case GOMP_MAP_ATTACH_DETACH:
case GOMP_MAP_DETACH:
break;
default:
map_seen |= 1;
@ -24200,7 +24209,9 @@ check_clauses:
case GOMP_MAP_PRESENT_ALLOC:
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_ALWAYS_POINTER:
case GOMP_MAP_POINTER:
case GOMP_MAP_ATTACH_DETACH:
case GOMP_MAP_ATTACH:
break;
default:
error_at (OMP_CLAUSE_LOCATION (*pc),

421
gcc/c/c-typeck.cc
View File

@ -13606,10 +13606,12 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
enum c_omp_region_type ort)
{
tree ret, low_bound, length, type;
bool openacc = (ort & C_ORT_ACC) != 0;
if (TREE_CODE (t) != TREE_LIST)
{
if (error_operand_p (t))
return error_mark_node;
c_omp_address_inspector ai (OMP_CLAUSE_LOCATION (c), t);
ret = t;
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_AFFINITY
&& OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND
@ -13619,59 +13621,17 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
t, omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
return error_mark_node;
}
while (INDIRECT_REF_P (t))
{
t = TREE_OPERAND (t, 0);
STRIP_NOPS (t);
if (TREE_CODE (t) == POINTER_PLUS_EXPR)
t = TREE_OPERAND (t, 0);
}
while (TREE_CODE (t) == COMPOUND_EXPR)
{
t = TREE_OPERAND (t, 1);
STRIP_NOPS (t);
}
if (TREE_CODE (t) == COMPONENT_REF
&& (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM))
{
if (DECL_BIT_FIELD (TREE_OPERAND (t, 1)))
{
error_at (OMP_CLAUSE_LOCATION (c),
"bit-field %qE in %qs clause",
t, omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
return error_mark_node;
}
while (TREE_CODE (t) == COMPONENT_REF)
{
if (TREE_CODE (TREE_TYPE (TREE_OPERAND (t, 0))) == UNION_TYPE)
{
error_at (OMP_CLAUSE_LOCATION (c),
"%qE is a member of a union", t);
return error_mark_node;
}
t = TREE_OPERAND (t, 0);
while (TREE_CODE (t) == MEM_REF
|| INDIRECT_REF_P (t)
|| TREE_CODE (t) == ARRAY_REF)
{
t = TREE_OPERAND (t, 0);
STRIP_NOPS (t);
if (TREE_CODE (t) == POINTER_PLUS_EXPR)
t = TREE_OPERAND (t, 0);
}
if (ort == C_ORT_ACC && TREE_CODE (t) == MEM_REF)
{
if (maybe_ne (mem_ref_offset (t), 0))
error_at (OMP_CLAUSE_LOCATION (c),
"cannot dereference %qE in %qs clause", t,
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
else
t = TREE_OPERAND (t, 0);
}
}
}
if (!ai.check_clause (c))
return error_mark_node;
else if (ai.component_access_p ()
&& (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM))
t = ai.get_root_term (true);
else
t = ai.unconverted_ref_origin ();
if (t == error_mark_node)
return error_mark_node;
if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
{
if (DECL_P (t))
@ -13766,7 +13726,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
{
error_at (OMP_CLAUSE_LOCATION (c),
"expected single pointer in %qs clause",
user_omp_clause_code_name (c, ort == C_ORT_ACC));
user_omp_clause_code_name (c, openacc));
return error_mark_node;
}
}
@ -13991,7 +13951,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
/* Handle array sections for clause C. */
static bool
handle_omp_array_sections (tree c, enum c_omp_region_type ort)
handle_omp_array_sections (tree &c, enum c_omp_region_type ort)
{
bool maybe_zero_len = false;
unsigned int first_non_one = 0;
@ -14200,58 +14160,47 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
OMP_CLAUSE_DECL (c) = first;
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
return false;
if (size)
size = c_fully_fold (size, false, NULL);
OMP_CLAUSE_SIZE (c) = size;
/* Don't set OMP_CLAUSE_SIZE for bare attach/detach clauses. */
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
|| (TREE_CODE (t) == COMPONENT_REF
&& TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE))
return false;
gcc_assert (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FORCE_DEVICEPTR);
switch (OMP_CLAUSE_MAP_KIND (c))
|| (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FORCE_DETACH))
{
case GOMP_MAP_ALLOC:
case GOMP_MAP_IF_PRESENT:
case GOMP_MAP_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_TOFROM:
case GOMP_MAP_ALWAYS_TO:
case GOMP_MAP_ALWAYS_FROM:
case GOMP_MAP_ALWAYS_TOFROM:
case GOMP_MAP_RELEASE:
case GOMP_MAP_DELETE:
case GOMP_MAP_FORCE_TO:
case GOMP_MAP_FORCE_FROM:
case GOMP_MAP_FORCE_TOFROM:
case GOMP_MAP_FORCE_PRESENT:
OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
break;
default:
break;
if (size)
size = c_fully_fold (size, false, NULL);
OMP_CLAUSE_SIZE (c) = size;
}
tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
if (TREE_CODE (t) == COMPONENT_REF)
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
else
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
OMP_CLAUSE_MAP_IMPLICIT (c2) = OMP_CLAUSE_MAP_IMPLICIT (c);
if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
&& !c_mark_addressable (t))
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
return false;
OMP_CLAUSE_DECL (c2) = t;
t = build_fold_addr_expr (first);
t = fold_convert_loc (OMP_CLAUSE_LOCATION (c), ptrdiff_type_node, t);
tree ptr = OMP_CLAUSE_DECL (c2);
if (!POINTER_TYPE_P (TREE_TYPE (ptr)))
ptr = build_fold_addr_expr (ptr);
t = fold_build2_loc (OMP_CLAUSE_LOCATION (c), MINUS_EXPR,
ptrdiff_type_node, t,
fold_convert_loc (OMP_CLAUSE_LOCATION (c),
ptrdiff_type_node, ptr));
t = c_fully_fold (t, false, NULL);
OMP_CLAUSE_SIZE (c2) = t;
OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
OMP_CLAUSE_CHAIN (c) = c2;
auto_vec<omp_addr_token *, 10> addr_tokens;
if (!omp_parse_expr (addr_tokens, first))
return true;
c_omp_address_inspector ai (OMP_CLAUSE_LOCATION (c), t);
tree nc = ai.expand_map_clause (c, first, addr_tokens, ort);
if (nc != error_mark_node)
{
using namespace omp_addr_tokenizer;
if (ai.maybe_zero_length_array_section (c))
OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
/* !!! If we're accessing a base decl via chained access
methods (e.g. multiple indirections), duplicate clause
detection won't work properly. Skip it in that case. */
if ((addr_tokens[0]->type == STRUCTURE_BASE
|| addr_tokens[0]->type == ARRAY_BASE)
&& addr_tokens[0]->u.structure_base_kind == BASE_DECL
&& addr_tokens[1]->type == ACCESS_METHOD
&& omp_access_chain_p (addr_tokens, 1))
c = nc;
return false;
}
}
return false;
}
@ -14517,7 +14466,6 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
tree ordered_clause = NULL_TREE;
tree schedule_clause = NULL_TREE;
bool oacc_async = false;
bool indir_component_ref_p = false;
tree last_iterators = NULL_TREE;
bool last_iterators_remove = false;
tree *nogroup_seen = NULL;
@ -14528,6 +14476,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
bool allocate_seen = false;
bool implicit_moved = false;
bool target_in_reduction_seen = false;
bool openacc = (ort & C_ORT_ACC) != 0;
bitmap_obstack_initialize (NULL);
bitmap_initialize (&generic_head, &bitmap_default_obstack);
@ -14543,7 +14492,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack);
bitmap_initialize (&is_on_device_head, &bitmap_default_obstack);
if (ort & C_ORT_ACC)
if (openacc)
for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_ASYNC)
{
@ -14937,8 +14886,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
else if ((ort == C_ORT_ACC
&& OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
else if ((openacc && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
|| (ort == C_ORT_OMP
&& (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR
|| (OMP_CLAUSE_CODE (c)
@ -14961,7 +14909,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
if (bitmap_bit_p (&oacc_reduction_head, DECL_UID (t)))
{
error_at (OMP_CLAUSE_LOCATION (c),
ort == C_ORT_ACC
openacc
? "%qD appears more than once in reduction clauses"
: "%qD appears more than once in data clauses",
t);
@ -14984,7 +14932,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR)
&& bitmap_bit_p (&map_head, DECL_UID (t)))
{
if (ort == C_ORT_ACC)
if (openacc)
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in data clauses", t);
else
@ -15049,9 +14997,10 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
"%qE appears more than once in data clauses", t);
remove = true;
}
else if (bitmap_bit_p (&map_head, DECL_UID (t)))
else if (bitmap_bit_p (&map_head, DECL_UID (t))
|| bitmap_bit_p (&map_field_head, DECL_UID (t)))
{
if (ort == C_ORT_ACC)
if (openacc)
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in data clauses", t);
else if (OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (c)
@ -15319,6 +15268,9 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
case OMP_CLAUSE_FROM:
case OMP_CLAUSE__CACHE_:
{
using namespace omp_addr_tokenizer;
auto_vec<omp_addr_token *, 10> addr_tokens;
t = OMP_CLAUSE_DECL (c);
if (TREE_CODE (t) == TREE_LIST)
{
@ -15347,56 +15299,68 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
}
while (TREE_CODE (t) == ARRAY_REF)
t = TREE_OPERAND (t, 0);
if (TREE_CODE (t) == COMPONENT_REF
&& TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE)
c_omp_address_inspector ai (OMP_CLAUSE_LOCATION (c), t);
if (!omp_parse_expr (addr_tokens, t))
{
do
{
t = TREE_OPERAND (t, 0);
if (TREE_CODE (t) == MEM_REF
|| INDIRECT_REF_P (t))
{
t = TREE_OPERAND (t, 0);
STRIP_NOPS (t);
if (TREE_CODE (t) == POINTER_PLUS_EXPR)
t = TREE_OPERAND (t, 0);
}
}
while (TREE_CODE (t) == COMPONENT_REF
|| TREE_CODE (t) == ARRAY_REF);
sorry_at (OMP_CLAUSE_LOCATION (c),
"unsupported map expression %qE",
OMP_CLAUSE_DECL (c));
remove = true;
break;
}
/* This check is to determine if this will be the only map
node created for this clause. Otherwise, we'll check
the following FIRSTPRIVATE_POINTER or ATTACH_DETACH
node on the next iteration(s) of the loop. */
if (addr_tokens.length () >= 4
&& addr_tokens[0]->type == STRUCTURE_BASE
&& addr_tokens[0]->u.structure_base_kind == BASE_DECL
&& addr_tokens[1]->type == ACCESS_METHOD
&& addr_tokens[2]->type == COMPONENT_SELECTOR
&& addr_tokens[3]->type == ACCESS_METHOD
&& (addr_tokens[3]->u.access_kind == ACCESS_DIRECT
|| (addr_tokens[3]->u.access_kind
== ACCESS_INDEXED_ARRAY)))
{
tree rt = addr_tokens[1]->expr;
gcc_assert (DECL_P (rt));
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_IMPLICIT (c)
&& (bitmap_bit_p (&map_head, DECL_UID (t))
|| bitmap_bit_p (&map_field_head, DECL_UID (t))
&& (bitmap_bit_p (&map_head, DECL_UID (rt))
|| bitmap_bit_p (&map_field_head, DECL_UID (rt))
|| bitmap_bit_p (&map_firstprivate_head,
DECL_UID (t))))
DECL_UID (rt))))
{
remove = true;
break;
}
if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
if (bitmap_bit_p (&map_field_head, DECL_UID (rt)))
break;
if (bitmap_bit_p (&map_head, DECL_UID (t)))
if (bitmap_bit_p (&map_head, DECL_UID (rt)))
{
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in motion "
"clauses", t);
else if (ort == C_ORT_ACC)
"clauses", rt);
else if (openacc)
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in data "
"clauses", t);
"clauses", rt);
else
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in map "
"clauses", t);
"clauses", rt);
remove = true;
}
else
{
bitmap_set_bit (&map_head, DECL_UID (t));
bitmap_set_bit (&map_field_head, DECL_UID (t));
bitmap_set_bit (&map_head, DECL_UID (rt));
bitmap_set_bit (&map_field_head, DECL_UID (rt));
}
}
}
@ -15404,7 +15368,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
remove = true;
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
&& !OMP_CLAUSE_SIZE (c))
/* In this case, we have a single array element which is a
pointer, and we already set OMP_CLAUSE_SIZE in
handle_omp_array_sections above. For attach/detach
@ -15413,6 +15378,14 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
OMP_CLAUSE_SIZE (c) = size_zero_node;
break;
}
else if (!omp_parse_expr (addr_tokens, t))
{
sorry_at (OMP_CLAUSE_LOCATION (c),
"unsupported map expression %qE",
OMP_CLAUSE_DECL (c));
remove = true;
break;
}
if (t == error_mark_node)
{
remove = true;
@ -15426,101 +15399,47 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
}
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
&& !OMP_CLAUSE_SIZE (c))
/* For attach/detach clauses, set OMP_CLAUSE_SIZE (representing a
bias) to zero here, so it is not set erroneously to the pointer
size later on in gimplify.cc. */
OMP_CLAUSE_SIZE (c) = size_zero_node;
while (INDIRECT_REF_P (t)
|| TREE_CODE (t) == ARRAY_REF)
c_omp_address_inspector ai (OMP_CLAUSE_LOCATION (c), t);
if (!ai.check_clause (c))
{
t = TREE_OPERAND (t, 0);
STRIP_NOPS (t);
if (TREE_CODE (t) == POINTER_PLUS_EXPR)
t = TREE_OPERAND (t, 0);
}
while (TREE_CODE (t) == COMPOUND_EXPR)
{
t = TREE_OPERAND (t, 1);
STRIP_NOPS (t);
}
indir_component_ref_p = false;
if (TREE_CODE (t) == COMPONENT_REF
&& (TREE_CODE (TREE_OPERAND (t, 0)) == MEM_REF
|| INDIRECT_REF_P (TREE_OPERAND (t, 0))
|| TREE_CODE (TREE_OPERAND (t, 0)) == ARRAY_REF))
{
t = TREE_OPERAND (TREE_OPERAND (t, 0), 0);
indir_component_ref_p = true;
STRIP_NOPS (t);
if (TREE_CODE (t) == POINTER_PLUS_EXPR)
t = TREE_OPERAND (t, 0);
remove = true;
break;
}
if (TREE_CODE (t) == COMPONENT_REF
&& OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_)
if (!ai.map_supported_p ())
{
if (DECL_BIT_FIELD (TREE_OPERAND (t, 1)))
{
error_at (OMP_CLAUSE_LOCATION (c),
"bit-field %qE in %qs clause",
t, omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
else if (!omp_mappable_type (TREE_TYPE (t)))
{
error_at (OMP_CLAUSE_LOCATION (c),
"%qE does not have a mappable type in %qs clause",
t, omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
else if (TYPE_ATOMIC (TREE_TYPE (t)))
{
error_at (OMP_CLAUSE_LOCATION (c),
"%<_Atomic%> %qE in %qs clause", t,
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
while (TREE_CODE (t) == COMPONENT_REF)
{
if (TREE_CODE (TREE_TYPE (TREE_OPERAND (t, 0)))
== UNION_TYPE)
{
error_at (OMP_CLAUSE_LOCATION (c),
"%qE is a member of a union", t);
remove = true;
break;
}
t = TREE_OPERAND (t, 0);
if (TREE_CODE (t) == MEM_REF)
{
if (maybe_ne (mem_ref_offset (t), 0))
error_at (OMP_CLAUSE_LOCATION (c),
"cannot dereference %qE in %qs clause", t,
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
else
t = TREE_OPERAND (t, 0);
}
while (TREE_CODE (t) == MEM_REF
|| INDIRECT_REF_P (t)
|| TREE_CODE (t) == ARRAY_REF)
{
t = TREE_OPERAND (t, 0);
STRIP_NOPS (t);
if (TREE_CODE (t) == POINTER_PLUS_EXPR)
t = TREE_OPERAND (t, 0);
}
}
if (remove)
break;
if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
{
if (bitmap_bit_p (&map_field_head, DECL_UID (t))
|| (ort != C_ORT_ACC
&& bitmap_bit_p (&map_head, DECL_UID (t))))
break;
}
sorry_at (OMP_CLAUSE_LOCATION (c),
"unsupported map expression %qE",
OMP_CLAUSE_DECL (c));
remove = true;
break;
}
gcc_assert ((addr_tokens[0]->type == ARRAY_BASE
|| addr_tokens[0]->type == STRUCTURE_BASE)
&& addr_tokens[1]->type == ACCESS_METHOD);
t = addr_tokens[1]->expr;
if (addr_tokens[0]->u.structure_base_kind != BASE_DECL)
goto skip_decl_checks;
/* For OpenMP, we can access a struct "t" and "t.d" on the same
mapping. OpenACC allows multiple fields of the same structure
to be written. */
if (addr_tokens[0]->type == STRUCTURE_BASE
&& (bitmap_bit_p (&map_field_head, DECL_UID (t))
|| (!openacc && bitmap_bit_p (&map_head, DECL_UID (t)))))
goto skip_decl_checks;
if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
{
error_at (OMP_CLAUSE_LOCATION (c),
@ -15538,7 +15457,6 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
else if ((OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
|| (OMP_CLAUSE_MAP_KIND (c)
!= GOMP_MAP_FIRSTPRIVATE_POINTER))
&& !indir_component_ref_p
&& !c_mark_addressable (t))
remove = true;
else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
@ -15584,27 +15502,25 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
remove = true;
}
else if (bitmap_bit_p (&map_head, DECL_UID (t))
&& !bitmap_bit_p (&map_field_head, DECL_UID (t)))
&& !bitmap_bit_p (&map_field_head, DECL_UID (t))
&& openacc)
{
if (ort == C_ORT_ACC)
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in data clauses",
t);
else
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears both in data and map clauses", t);
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in data clauses", t);
remove = true;
}
else
bitmap_set_bit (&map_firstprivate_head, DECL_UID (t));
}
else if (bitmap_bit_p (&map_head, DECL_UID (t))
&& !bitmap_bit_p (&map_field_head, DECL_UID (t)))
&& !bitmap_bit_p (&map_field_head, DECL_UID (t))
&& ort != C_ORT_OMP
&& ort != C_ORT_OMP_EXIT_DATA)
{
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in motion clauses", t);
else if (ort == C_ORT_ACC)
else if (openacc)
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in data clauses", t);
else
@ -15612,8 +15528,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
"%qD appears more than once in map clauses", t);
remove = true;
}
else if (ort == C_ORT_ACC
&& bitmap_bit_p (&generic_head, DECL_UID (t)))
else if (openacc && bitmap_bit_p (&generic_head, DECL_UID (t)))
{
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in data clauses", t);
@ -15622,7 +15537,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
else if (bitmap_bit_p (&firstprivate_head, DECL_UID (t))
|| bitmap_bit_p (&is_on_device_head, DECL_UID (t)))
{
if (ort == C_ORT_ACC)
if (openacc)
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in data clauses", t);
else
@ -15630,13 +15545,37 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
"%qD appears both in data and map clauses", t);
remove = true;
}
else
else if (!omp_access_chain_p (addr_tokens, 1))
{
bitmap_set_bit (&map_head, DECL_UID (t));
if (t != OMP_CLAUSE_DECL (c)
&& TREE_CODE (OMP_CLAUSE_DECL (c)) == COMPONENT_REF)
bitmap_set_bit (&map_field_head, DECL_UID (t));
}
skip_decl_checks:
/* If we call omp_expand_map_clause in handle_omp_array_sections,
the containing loop (here) iterates through the new nodes
created by that expansion. Avoid expanding those again (just
by checking the node type). */
if (!remove
&& ort != C_ORT_DECLARE_SIMD
&& (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
|| (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER
&& (OMP_CLAUSE_MAP_KIND (c)
!= GOMP_MAP_FIRSTPRIVATE_REFERENCE)
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH)))
{
grp_start_p = pc;
grp_sentinel = OMP_CLAUSE_CHAIN (c);
tree nc = ai.expand_map_clause (c, OMP_CLAUSE_DECL (c),
addr_tokens, ort);
if (nc != error_mark_node)
c = nc;
}
}
break;
@ -15714,7 +15653,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
if (TREE_CODE (TREE_TYPE (t)) != POINTER_TYPE)
{
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR
&& ort != C_ORT_ACC)
&& !openacc)
{
error_at (OMP_CLAUSE_LOCATION (c),
"%qs variable is not a pointer",

23
gcc/cp/parser.cc
View File

@ -41426,7 +41426,7 @@ cp_parser_oacc_compute_clause_self (cp_parser *parser, tree list)
static tree
cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
const char *where, cp_token *pragma_tok,
bool finish_p = true)
bool finish_p = true, bool target_p = false)
{
tree clauses = NULL;
bool first = true;
@ -41645,7 +41645,8 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
cp_parser_skip_to_pragma_eol (parser, pragma_tok);
if (finish_p)
return finish_omp_clauses (clauses, C_ORT_ACC);
return finish_omp_clauses (clauses, target_p ? C_ORT_ACC_TARGET
: C_ORT_ACC);
return clauses;
}
@ -46236,6 +46237,7 @@ cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p)
case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
case GOMP_MAP_ALWAYS_POINTER:
case GOMP_MAP_ATTACH_DETACH:
case GOMP_MAP_ATTACH:
break;
default:
map_seen |= 1;
@ -46358,6 +46360,7 @@ cp_parser_omp_target_enter_data (cp_parser *parser, cp_token *pragma_tok,
case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
case GOMP_MAP_ALWAYS_POINTER:
case GOMP_MAP_ATTACH_DETACH:
case GOMP_MAP_ATTACH:
break;
default:
map_seen |= 1;
@ -46436,7 +46439,9 @@ cp_parser_omp_target_exit_data (cp_parser *parser, cp_token *pragma_tok,
tree clauses
= cp_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK,
"#pragma omp target exit data", pragma_tok);
"#pragma omp target exit data", pragma_tok,
false);
clauses = finish_omp_clauses (clauses, C_ORT_OMP_EXIT_DATA);
c_omp_adjust_map_clauses (clauses, false);
int map_seen = 0;
for (tree *pc = &clauses; *pc;)
@ -46472,6 +46477,7 @@ cp_parser_omp_target_exit_data (cp_parser *parser, cp_token *pragma_tok,
case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
case GOMP_MAP_ALWAYS_POINTER:
case GOMP_MAP_ATTACH_DETACH:
case GOMP_MAP_DETACH:
break;
default:
map_seen |= 1;
@ -46853,7 +46859,7 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
bool found_in_scope = global_bindings_p ();
clauses = cp_parser_oacc_all_clauses (parser, OACC_DECLARE_CLAUSE_MASK,
"#pragma acc declare", pragma_tok, true);
"#pragma acc declare", pragma_tok);
if (omp_find_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE)
@ -47101,12 +47107,13 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
mask |= OACC_LOOP_CLAUSE_MASK;
tree clauses = cp_parser_oacc_all_clauses (parser, mask, p_name, pragma_tok,
cclauses == NULL);
/*finish_p=*/cclauses == NULL,
/*target=*/is_parallel);
if (cclauses)
{
clauses = c_oacc_split_loop_clauses (clauses, cclauses, is_parallel);
if (*cclauses)
*cclauses = finish_omp_clauses (*cclauses, C_ORT_ACC);
*cclauses = finish_omp_clauses (*cclauses, C_ORT_ACC_TARGET);
if (clauses)
clauses = finish_omp_clauses (clauses, C_ORT_ACC);
}
@ -47242,7 +47249,9 @@ cp_parser_oacc_compute (cp_parser *parser, cp_token *pragma_tok,
}
}
tree clauses = cp_parser_oacc_all_clauses (parser, mask, p_name, pragma_tok);
tree clauses = cp_parser_oacc_all_clauses (parser, mask, p_name, pragma_tok,
/*finish_p=*/true,
/*target=*/true);
tree block = begin_omp_parallel ();
unsigned int save = cp_parser_begin_omp_structured_block (parser);

4
gcc/cp/pt.cc
View File

@ -18727,8 +18727,8 @@ tsubst_stmt (tree t, tree args, tsubst_flags_t complain, tree in_decl)
case OACC_KERNELS:
case OACC_PARALLEL:
case OACC_SERIAL:
tmp = tsubst_omp_clauses (OMP_CLAUSES (t), C_ORT_ACC, args, complain,
in_decl);
tmp = tsubst_omp_clauses (OMP_CLAUSES (t), C_ORT_ACC_TARGET, args,
complain, in_decl);
stmt = begin_omp_parallel ();
RECUR (OMP_BODY (t));
finish_omp_construct (TREE_CODE (t), stmt, tmp);

633
gcc/cp/semantics.cc
View File

@ -5273,6 +5273,54 @@ omp_privatize_field (tree t, bool shared)
return v;
}
/* C++ specialisation of the c_omp_address_inspector class. */
class cp_omp_address_inspector : public c_omp_address_inspector
{
public:
cp_omp_address_inspector (location_t loc, tree t)
: c_omp_address_inspector (loc, t)
{
}
~cp_omp_address_inspector ()
{
}
bool processing_template_decl_p ()
{
return processing_template_decl;
}
void emit_unmappable_type_notes (tree t)
{
if (TREE_TYPE (t) != error_mark_node
&& !COMPLETE_TYPE_P (TREE_TYPE (t)))
cxx_incomplete_type_inform (TREE_TYPE (t));
}
tree convert_from_reference (tree x)
{
return ::convert_from_reference (x);
}
tree build_array_ref (location_t loc, tree arr, tree idx)
{
return ::build_array_ref (loc, arr, idx);
}
bool check_clause (tree clause)
{
if (TREE_CODE (orig) == COMPONENT_REF
&& invalid_nonstatic_memfn_p (EXPR_LOCATION (orig), orig,
tf_warning_or_error))
return false;
if (!c_omp_address_inspector::check_clause (clause))
return false;
return true;
}
};
/* Helper function for handle_omp_array_sections. Called recursively
to handle multiple array-section-subscripts. C is the clause,
T current expression (initially OMP_CLAUSE_DECL), which is either
@ -5299,63 +5347,27 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
enum c_omp_region_type ort)
{
tree ret, low_bound, length, type;
bool openacc = (ort & C_ORT_ACC) != 0;
if (TREE_CODE (t) != TREE_LIST)
{
if (error_operand_p (t))
return error_mark_node;
if (REFERENCE_REF_P (t)
&& TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
t = TREE_OPERAND (t, 0);
ret = t;
while (INDIRECT_REF_P (t))
{
t = TREE_OPERAND (t, 0);
STRIP_NOPS (t);
if (TREE_CODE (t) == POINTER_PLUS_EXPR)
t = TREE_OPERAND (t, 0);
}
while (TREE_CODE (t) == COMPOUND_EXPR)
{
t = TREE_OPERAND (t, 1);
STRIP_NOPS (t);
}
if (TREE_CODE (t) == COMPONENT_REF
&& (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM)
&& !type_dependent_expression_p (t))
{
if (TREE_CODE (TREE_OPERAND (t, 1)) == FIELD_DECL
&& DECL_BIT_FIELD (TREE_OPERAND (t, 1)))
{
error_at (OMP_CLAUSE_LOCATION (c),
"bit-field %qE in %qs clause",
t, omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
return error_mark_node;
}
while (TREE_CODE (t) == COMPONENT_REF)
{
if (TREE_TYPE (TREE_OPERAND (t, 0))
&& TREE_CODE (TREE_TYPE (TREE_OPERAND (t, 0))) == UNION_TYPE)
{
error_at (OMP_CLAUSE_LOCATION (c),
"%qE is a member of a union", t);
return error_mark_node;
}
t = TREE_OPERAND (t, 0);
while (TREE_CODE (t) == MEM_REF
|| INDIRECT_REF_P (t)
|| TREE_CODE (t) == ARRAY_REF)
{
t = TREE_OPERAND (t, 0);
STRIP_NOPS (t);
if (TREE_CODE (t) == POINTER_PLUS_EXPR)
t = TREE_OPERAND (t, 0);
}
}
if (REFERENCE_REF_P (t))
t = TREE_OPERAND (t, 0);
}
cp_omp_address_inspector ai (OMP_CLAUSE_LOCATION (c), t);
tree t_refto = ai.maybe_unconvert_ref (t);
if (!ai.check_clause (c))
return error_mark_node;
else if (ai.component_access_p ()
&& (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM))
t = ai.get_root_term (true);
else
t = ai.unconverted_ref_origin ();
if (t == error_mark_node)
return error_mark_node;
ret = t_refto;
if (TREE_CODE (t) == FIELD_DECL)
ret = finish_non_static_data_member (t, NULL_TREE, NULL_TREE);
else if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
@ -5452,7 +5464,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
{
error_at (OMP_CLAUSE_LOCATION (c),
"expected single pointer in %qs clause",
user_omp_clause_code_name (c, ort == C_ORT_ACC));
user_omp_clause_code_name (c, openacc));
return error_mark_node;
}
}
@ -5689,7 +5701,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
/* Handle array sections for clause C. */
static bool
handle_omp_array_sections (tree c, enum c_omp_region_type ort)
handle_omp_array_sections (tree &c, enum c_omp_region_type ort)
{
bool maybe_zero_len = false;
unsigned int first_non_one = 0;
@ -5897,118 +5909,73 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
OMP_CLAUSE_DECL (c) = first;
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
return false;
OMP_CLAUSE_SIZE (c) = size;
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
|| (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FORCE_DETACH))
OMP_CLAUSE_SIZE (c) = size;
if (TREE_CODE (t) == FIELD_DECL)
t = finish_non_static_data_member (t, NULL_TREE, NULL_TREE);
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
|| (TREE_CODE (t) == COMPONENT_REF
&& TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE))
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
return false;
switch (OMP_CLAUSE_MAP_KIND (c))
{
case GOMP_MAP_ALLOC:
case GOMP_MAP_IF_PRESENT:
case GOMP_MAP_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_TOFROM:
case GOMP_MAP_ALWAYS_TO:
case GOMP_MAP_ALWAYS_FROM:
case GOMP_MAP_ALWAYS_TOFROM:
case GOMP_MAP_PRESENT_ALLOC:
case GOMP_MAP_PRESENT_TO:
case GOMP_MAP_PRESENT_FROM:
case GOMP_MAP_PRESENT_TOFROM:
case GOMP_MAP_ALWAYS_PRESENT_TO:
case GOMP_MAP_ALWAYS_PRESENT_FROM:
case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
case GOMP_MAP_RELEASE:
case GOMP_MAP_DELETE:
case GOMP_MAP_FORCE_TO:
case GOMP_MAP_FORCE_FROM:
case GOMP_MAP_FORCE_TOFROM:
case GOMP_MAP_FORCE_PRESENT:
OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
break;
default:
break;
}
bool reference_always_pointer = true;
tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
OMP_CLAUSE_MAP);
if (TREE_CODE (t) == COMPONENT_REF)
{
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
if ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP
&& TYPE_REF_P (TREE_TYPE (t)))
if (TREE_CODE (first) == INDIRECT_REF)
{
/* Detect and skip adding extra nodes for pointer-to-member
mappings. These are unsupported for now. */
tree tmp = TREE_OPERAND (first, 0);
if (TREE_CODE (tmp) == NON_LVALUE_EXPR)
tmp = TREE_OPERAND (tmp, 0);
if (TREE_CODE (tmp) == INDIRECT_REF)
tmp = TREE_OPERAND (tmp, 0);
if (TREE_CODE (tmp) == POINTER_PLUS_EXPR)
{
if (TREE_CODE (TREE_TYPE (TREE_TYPE (t))) == ARRAY_TYPE)
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
else
t = convert_from_reference (t);
reference_always_pointer = false;
tree offset = TREE_OPERAND (tmp, 1);
STRIP_NOPS (offset);
if (TYPE_PTRMEM_P (TREE_TYPE (offset)))
{
sorry_at (OMP_CLAUSE_LOCATION (c),
"pointer-to-member mapping %qE not supported",
OMP_CLAUSE_DECL (c));
return true;
}
}
}
else if (REFERENCE_REF_P (t)
&& TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
{
gomp_map_kind k;
if ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP
&& TREE_CODE (TREE_TYPE (t)) == POINTER_TYPE)
k = GOMP_MAP_ATTACH_DETACH;
else
{
t = TREE_OPERAND (t, 0);
k = (ort == C_ORT_ACC
? GOMP_MAP_ATTACH_DETACH : GOMP_MAP_ALWAYS_POINTER);
}
OMP_CLAUSE_SET_MAP_KIND (c2, k);
}
else
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
OMP_CLAUSE_MAP_IMPLICIT (c2) = OMP_CLAUSE_MAP_IMPLICIT (c);
if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
&& !cxx_mark_addressable (t))
return false;
OMP_CLAUSE_DECL (c2) = t;
t = build_fold_addr_expr (first);
t = fold_convert_loc (OMP_CLAUSE_LOCATION (c),
ptrdiff_type_node, t);
tree ptr = OMP_CLAUSE_DECL (c2);
ptr = convert_from_reference (ptr);
if (!INDIRECT_TYPE_P (TREE_TYPE (ptr)))
ptr = build_fold_addr_expr (ptr);
t = fold_build2_loc (OMP_CLAUSE_LOCATION (c), MINUS_EXPR,
ptrdiff_type_node, t,
fold_convert_loc (OMP_CLAUSE_LOCATION (c),
ptrdiff_type_node, ptr));
OMP_CLAUSE_SIZE (c2) = t;
OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
OMP_CLAUSE_CHAIN (c) = c2;
ptr = OMP_CLAUSE_DECL (c2);
if (reference_always_pointer
&& OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
&& TYPE_REF_P (TREE_TYPE (ptr))
&& INDIRECT_TYPE_P (TREE_TYPE (TREE_TYPE (ptr))))
/* FIRST represents the first item of data that we are mapping.
E.g. if we're mapping an array, FIRST might resemble
"foo.bar.myarray[0]". */
auto_vec<omp_addr_token *, 10> addr_tokens;
if (!omp_parse_expr (addr_tokens, first))
return true;
cp_omp_address_inspector ai (OMP_CLAUSE_LOCATION (c), t);
tree nc = ai.expand_map_clause (c, first, addr_tokens, ort);
if (nc != error_mark_node)
{
tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (c3, OMP_CLAUSE_MAP_KIND (c2));
OMP_CLAUSE_MAP_IMPLICIT (c2) = OMP_CLAUSE_MAP_IMPLICIT (c);
OMP_CLAUSE_DECL (c3) = ptr;
if (OMP_CLAUSE_MAP_KIND (c2) == GOMP_MAP_ALWAYS_POINTER
|| OMP_CLAUSE_MAP_KIND (c2) == GOMP_MAP_ATTACH_DETACH)
{
OMP_CLAUSE_DECL (c2) = build_simple_mem_ref (ptr);
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
}
else
OMP_CLAUSE_DECL (c2) = convert_from_reference (ptr);
OMP_CLAUSE_SIZE (c3) = size_zero_node;
OMP_CLAUSE_CHAIN (c3) = OMP_CLAUSE_CHAIN (c2);
OMP_CLAUSE_CHAIN (c2) = c3;
using namespace omp_addr_tokenizer;
if (ai.maybe_zero_length_array_section (c))
OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
/* !!! If we're accessing a base decl via chained access
methods (e.g. multiple indirections), duplicate clause
detection won't work properly. Skip it in that case. */
if ((addr_tokens[0]->type == STRUCTURE_BASE
|| addr_tokens[0]->type == ARRAY_BASE)
&& addr_tokens[0]->u.structure_base_kind == BASE_DECL
&& addr_tokens[1]->type == ACCESS_METHOD
&& omp_access_chain_p (addr_tokens, 1))
c = nc;
return false;
}
}
}
@ -6936,6 +6903,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
bitmap_head oacc_reduction_head, is_on_device_head;
tree c, t, *pc;
tree safelen = NULL_TREE;
bool openacc = (ort & C_ORT_ACC) != 0;
bool branch_seen = false;
bool copyprivate_seen = false;
bool ordered_seen = false;
@ -6968,7 +6936,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack);
bitmap_initialize (&is_on_device_head, &bitmap_default_obstack);
if (ort & C_ORT_ACC)
if (openacc)
for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_ASYNC)
{
@ -7217,7 +7185,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
t = OMP_CLAUSE_DECL (c);
check_dup_generic_t:
if (t == current_class_ptr
&& ((ort != C_ORT_OMP_DECLARE_SIMD && ort != C_ORT_ACC)
&& ((ort != C_ORT_OMP_DECLARE_SIMD && !openacc)
|| (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_LINEAR
&& OMP_CLAUSE_CODE (c) != OMP_CLAUSE_UNIFORM)))
{
@ -7242,7 +7210,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
else if ((ort == C_ORT_ACC
else if ((openacc
&& OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
|| (ort == C_ORT_OMP
&& (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR
@ -7266,7 +7234,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
if (bitmap_bit_p (&oacc_reduction_head, DECL_UID (t)))
{
error_at (OMP_CLAUSE_LOCATION (c),
ort == C_ORT_ACC
openacc
? "%qD appears more than once in reduction clauses"
: "%qD appears more than once in data clauses",
t);
@ -7289,7 +7257,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR)
&& bitmap_bit_p (&map_head, DECL_UID (t)))
{
if (ort == C_ORT_ACC)
if (openacc)
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in data clauses", t);
else
@ -7351,7 +7319,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
omp_note_field_privatization (t, OMP_CLAUSE_DECL (c));
else
t = OMP_CLAUSE_DECL (c);
if (ort != C_ORT_ACC && t == current_class_ptr)
if (!openacc && t == current_class_ptr)
{
error_at (OMP_CLAUSE_LOCATION (c),
"%<this%> allowed in OpenMP only in %<declare simd%>"
@ -7387,9 +7355,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
"%qD appears more than once in data clauses", t);
remove = true;
}
else if (bitmap_bit_p (&map_head, DECL_UID (t)))
else if (bitmap_bit_p (&map_head, DECL_UID (t))
|| bitmap_bit_p (&map_field_head, DECL_UID (t)))
{
if (ort == C_ORT_ACC)
if (openacc)
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in data clauses", t);
else if (OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (c)
@ -7410,7 +7379,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
omp_note_field_privatization (t, OMP_CLAUSE_DECL (c));
else
t = OMP_CLAUSE_DECL (c);
if (ort != C_ORT_ACC && t == current_class_ptr)
if (!openacc && t == current_class_ptr)
{
error_at (OMP_CLAUSE_LOCATION (c),
"%<this%> allowed in OpenMP only in %<declare simd%>"
@ -8211,6 +8180,9 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
case OMP_CLAUSE_FROM:
case OMP_CLAUSE__CACHE_:
{
using namespace omp_addr_tokenizer;
auto_vec<omp_addr_token *, 10> addr_tokens;
t = OMP_CLAUSE_DECL (c);
if (TREE_CODE (t) == TREE_LIST)
{
@ -8237,58 +8209,73 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
}
while (TREE_CODE (t) == ARRAY_REF)
t = TREE_OPERAND (t, 0);
if (TREE_CODE (t) == COMPONENT_REF
&& TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE)
if (type_dependent_expression_p (t))
break;
cp_omp_address_inspector ai (OMP_CLAUSE_LOCATION (c), t);
if (!ai.map_supported_p ()
|| !omp_parse_expr (addr_tokens, t))
{
do
{
t = TREE_OPERAND (t, 0);
if (REFERENCE_REF_P (t))
t = TREE_OPERAND (t, 0);
if (TREE_CODE (t) == MEM_REF
|| INDIRECT_REF_P (t))
{
t = TREE_OPERAND (t, 0);
STRIP_NOPS (t);
if (TREE_CODE (t) == POINTER_PLUS_EXPR)
t = TREE_OPERAND (t, 0);
}
}
while (TREE_CODE (t) == COMPONENT_REF
|| TREE_CODE (t) == ARRAY_REF);
sorry_at (OMP_CLAUSE_LOCATION (c),
"unsupported map expression %qE",
OMP_CLAUSE_DECL (c));
remove = true;
break;
}
/* This check is to determine if this will be the only map
node created for this clause. Otherwise, we'll check
the following FIRSTPRIVATE_POINTER,
FIRSTPRIVATE_REFERENCE or ATTACH_DETACH node on the next
iteration(s) of the loop. */
if (addr_tokens.length () >= 4
&& addr_tokens[0]->type == STRUCTURE_BASE
&& addr_tokens[0]->u.structure_base_kind == BASE_DECL
&& addr_tokens[1]->type == ACCESS_METHOD
&& addr_tokens[2]->type == COMPONENT_SELECTOR
&& addr_tokens[3]->type == ACCESS_METHOD
&& (addr_tokens[3]->u.access_kind == ACCESS_DIRECT
|| (addr_tokens[3]->u.access_kind
== ACCESS_INDEXED_ARRAY)))
{
tree rt = addr_tokens[1]->expr;
gcc_assert (DECL_P (rt));
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_IMPLICIT (c)
&& (bitmap_bit_p (&map_head, DECL_UID (t))
|| bitmap_bit_p (&map_field_head, DECL_UID (t))
&& (bitmap_bit_p (&map_head, DECL_UID (rt))
|| bitmap_bit_p (&map_field_head, DECL_UID (rt))
|| bitmap_bit_p (&map_firstprivate_head,
DECL_UID (t))))
DECL_UID (rt))))
{
remove = true;
break;
}
if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
if (bitmap_bit_p (&map_field_head, DECL_UID (rt)))
break;
if (bitmap_bit_p (&map_head, DECL_UID (t)))
if (bitmap_bit_p (&map_head, DECL_UID (rt)))
{
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in motion"
" clauses", t);
else if (ort == C_ORT_ACC)
" clauses", rt);
else if (openacc)
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in data"
" clauses", t);
" clauses", rt);
else
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in map"
" clauses", t);
" clauses", rt);
remove = true;
}
else
{
bitmap_set_bit (&map_head, DECL_UID (t));
bitmap_set_bit (&map_field_head, DECL_UID (t));
bitmap_set_bit (&map_head, DECL_UID (rt));
bitmap_set_bit (&map_field_head, DECL_UID (rt));
}
}
}
@ -8296,7 +8283,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
remove = true;
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
&& !OMP_CLAUSE_SIZE (c))
/* In this case, we have a single array element which is a
pointer, and we already set OMP_CLAUSE_SIZE in
handle_omp_array_sections above. For attach/detach
@ -8305,6 +8293,16 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
OMP_CLAUSE_SIZE (c) = size_zero_node;
break;
}
else if (type_dependent_expression_p (t))
break;
else if (!omp_parse_expr (addr_tokens, t))
{
sorry_at (OMP_CLAUSE_LOCATION (c),
"unsupported map expression %qE",
OMP_CLAUSE_DECL (c));
remove = true;
break;
}
if (t == error_mark_node)
{
remove = true;
@ -8318,115 +8316,55 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
}
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
&& !OMP_CLAUSE_SIZE (c))
/* For attach/detach clauses, set OMP_CLAUSE_SIZE (representing a
bias) to zero here, so it is not set erroneously to the
pointer size later on in gimplify.cc. */
OMP_CLAUSE_SIZE (c) = size_zero_node;
if (REFERENCE_REF_P (t)
&& TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
cp_omp_address_inspector ai (OMP_CLAUSE_LOCATION (c), t);
if (!ai.check_clause (c))
{
t = TREE_OPERAND (t, 0);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH)
OMP_CLAUSE_DECL (c) = t;
remove = true;
break;
}
while (INDIRECT_REF_P (t)
|| TREE_CODE (t) == ARRAY_REF)
if (!ai.map_supported_p ())
{
t = TREE_OPERAND (t, 0);
STRIP_NOPS (t);
if (TREE_CODE (t) == POINTER_PLUS_EXPR)
t = TREE_OPERAND (t, 0);
sorry_at (OMP_CLAUSE_LOCATION (c),
"unsupported map expression %qE",
OMP_CLAUSE_DECL (c));
remove = true;
break;
}
while (TREE_CODE (t) == COMPOUND_EXPR)
{
t = TREE_OPERAND (t, 1);
STRIP_NOPS (t);
}
if (TREE_CODE (t) == COMPONENT_REF
&& invalid_nonstatic_memfn_p (EXPR_LOCATION (t), t,
tf_warning_or_error))
remove = true;
indir_component_ref_p = false;
if (TREE_CODE (t) == COMPONENT_REF
&& (INDIRECT_REF_P (TREE_OPERAND (t, 0))
|| TREE_CODE (TREE_OPERAND (t, 0)) == ARRAY_REF))
{
t = TREE_OPERAND (TREE_OPERAND (t, 0), 0);
indir_component_ref_p = true;
STRIP_NOPS (t);
if (TREE_CODE (t) == POINTER_PLUS_EXPR)
t = TREE_OPERAND (t, 0);
}
if (TREE_CODE (t) == COMPONENT_REF
&& OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_)
{
if (type_dependent_expression_p (t))
break;
if (TREE_CODE (TREE_OPERAND (t, 1)) == FIELD_DECL
&& DECL_BIT_FIELD (TREE_OPERAND (t, 1)))
{
error_at (OMP_CLAUSE_LOCATION (c),
"bit-field %qE in %qs clause",
t, omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
else if (!omp_mappable_type (TREE_TYPE (t)))
{
error_at (OMP_CLAUSE_LOCATION (c),
"%qE does not have a mappable type in %qs clause",
t, omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
if (TREE_TYPE (t) != error_mark_node
&& !COMPLETE_TYPE_P (TREE_TYPE (t)))
cxx_incomplete_type_inform (TREE_TYPE (t));
remove = true;
}
while (TREE_CODE (t) == COMPONENT_REF)
{
if (TREE_TYPE (TREE_OPERAND (t, 0))
&& (TREE_CODE (TREE_TYPE (TREE_OPERAND (t, 0)))
== UNION_TYPE))
{
error_at (OMP_CLAUSE_LOCATION (c),
"%qE is a member of a union", t);
remove = true;
break;
}
t = TREE_OPERAND (t, 0);
if (TREE_CODE (t) == MEM_REF)
{
if (maybe_ne (mem_ref_offset (t), 0))
error_at (OMP_CLAUSE_LOCATION (c),
"cannot dereference %qE in %qs clause", t,
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
else
t = TREE_OPERAND (t, 0);
}
while (TREE_CODE (t) == MEM_REF
|| INDIRECT_REF_P (t)
|| TREE_CODE (t) == ARRAY_REF)
{
t = TREE_OPERAND (t, 0);
STRIP_NOPS (t);
if (TREE_CODE (t) == POINTER_PLUS_EXPR)
t = TREE_OPERAND (t, 0);
}
}
if (remove)
break;
if (REFERENCE_REF_P (t))
t = TREE_OPERAND (t, 0);
if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
{
if (bitmap_bit_p (&map_field_head, DECL_UID (t))
|| (ort != C_ORT_ACC
&& bitmap_bit_p (&map_head, DECL_UID (t))))
goto handle_map_references;
}
}
if (!processing_template_decl
&& TREE_CODE (t) == FIELD_DECL)
gcc_assert ((addr_tokens[0]->type == ARRAY_BASE
|| addr_tokens[0]->type == STRUCTURE_BASE)
&& addr_tokens[1]->type == ACCESS_METHOD);
t = addr_tokens[1]->expr;
/* This is used to prevent cxx_mark_addressable from being called
on 'this' for expressions like 'this->a', i.e. typical member
accesses. */
indir_component_ref_p
= (addr_tokens[0]->type == STRUCTURE_BASE
&& addr_tokens[1]->u.access_kind != ACCESS_DIRECT);
if (addr_tokens[0]->u.structure_base_kind != BASE_DECL)
goto skip_decl_checks;
/* For OpenMP, we can access a struct "t" and "t.d" on the same
mapping. OpenACC allows multiple fields of the same structure
to be written. */
if (addr_tokens[0]->type == STRUCTURE_BASE
&& (bitmap_bit_p (&map_field_head, DECL_UID (t))
|| (!openacc && bitmap_bit_p (&map_head, DECL_UID (t)))))
goto skip_decl_checks;
if (!processing_template_decl && TREE_CODE (t) == FIELD_DECL)
{
OMP_CLAUSE_DECL (c)
= finish_non_static_data_member (t, NULL_TREE, NULL_TREE);
@ -8464,12 +8402,17 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
|| (OMP_CLAUSE_MAP_KIND (c)
!= GOMP_MAP_FIRSTPRIVATE_POINTER))
&& !indir_component_ref_p
&& (t != current_class_ptr
|| OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
|| OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH)
&& !cxx_mark_addressable (t))
remove = true;
else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
|| (OMP_CLAUSE_MAP_KIND (c)
== GOMP_MAP_FIRSTPRIVATE_POINTER)))
== GOMP_MAP_FIRSTPRIVATE_POINTER)
|| (OMP_CLAUSE_MAP_KIND (c)
== GOMP_MAP_ATTACH_DETACH)))
&& t == OMP_CLAUSE_DECL (c)
&& !type_dependent_expression_p (t)
&& !omp_mappable_type (TYPE_REF_P (TREE_TYPE (t))
@ -8513,27 +8456,29 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
remove = true;
}
else if (bitmap_bit_p (&map_head, DECL_UID (t))
&& !bitmap_bit_p (&map_field_head, DECL_UID (t)))
&& !bitmap_bit_p (&map_field_head, DECL_UID (t))
&& openacc)
{
if (ort == C_ORT_ACC)
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in data clauses",
t);
else
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears both in data and map clauses", t);
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in data clauses", t);
remove = true;
}
else
bitmap_set_bit (&map_firstprivate_head, DECL_UID (t));
}
else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& (OMP_CLAUSE_MAP_KIND (c)
== GOMP_MAP_FIRSTPRIVATE_REFERENCE))
bitmap_set_bit (&map_firstprivate_head, DECL_UID (t));
else if (bitmap_bit_p (&map_head, DECL_UID (t))
&& !bitmap_bit_p (&map_field_head, DECL_UID (t)))
&& !bitmap_bit_p (&map_field_head, DECL_UID (t))
&& ort != C_ORT_OMP
&& ort != C_ORT_OMP_EXIT_DATA)
{
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in motion clauses", t);
else if (ort == C_ORT_ACC)
else if (openacc)
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in data clauses", t);
else
@ -8541,8 +8486,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
"%qD appears more than once in map clauses", t);
remove = true;
}
else if (ort == C_ORT_ACC
&& bitmap_bit_p (&generic_head, DECL_UID (t)))
else if (openacc && bitmap_bit_p (&generic_head, DECL_UID (t)))
{
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in data clauses", t);
@ -8551,7 +8495,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
else if (bitmap_bit_p (&firstprivate_head, DECL_UID (t))
|| bitmap_bit_p (&is_on_device_head, DECL_UID (t)))
{
if (ort == C_ORT_ACC)
if (openacc)
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in data clauses", t);
else
@ -8559,7 +8503,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
"%qD appears both in data and map clauses", t);
remove = true;
}
else
else if (!omp_access_chain_p (addr_tokens, 1))
{
bitmap_set_bit (&map_head, DECL_UID (t));
@ -8573,49 +8517,30 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
0))))))
bitmap_set_bit (&map_field_head, DECL_UID (t));
}
handle_map_references:
skip_decl_checks:
/* If we call ai.expand_map_clause in handle_omp_array_sections,
the containing loop (here) iterates through the new nodes
created by that expansion. Avoid expanding those again (just
by checking the node type). */
if (!remove
&& !processing_template_decl
&& ort != C_ORT_DECLARE_SIMD
&& TYPE_REF_P (TREE_TYPE (OMP_CLAUSE_DECL (c))))
&& (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
|| (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER
&& (OMP_CLAUSE_MAP_KIND (c)
!= GOMP_MAP_FIRSTPRIVATE_REFERENCE)
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH)))
{
t = OMP_CLAUSE_DECL (c);
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
{
OMP_CLAUSE_DECL (c) = build_simple_mem_ref (t);
if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
OMP_CLAUSE_SIZE (c)
= TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (t)));
}
else if (OMP_CLAUSE_MAP_KIND (c)
!= GOMP_MAP_FIRSTPRIVATE_POINTER
&& (OMP_CLAUSE_MAP_KIND (c)
!= GOMP_MAP_FIRSTPRIVATE_REFERENCE)
&& (OMP_CLAUSE_MAP_KIND (c)
!= GOMP_MAP_ALWAYS_POINTER)
&& (OMP_CLAUSE_MAP_KIND (c)
!= GOMP_MAP_ATTACH_DETACH))
{
grp_start_p = pc;
grp_sentinel = OMP_CLAUSE_CHAIN (c);
tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
OMP_CLAUSE_MAP);
if (TREE_CODE (t) == COMPONENT_REF)
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
else
OMP_CLAUSE_SET_MAP_KIND (c2,
GOMP_MAP_FIRSTPRIVATE_REFERENCE);
OMP_CLAUSE_DECL (c2) = t;
OMP_CLAUSE_SIZE (c2) = size_zero_node;
OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
OMP_CLAUSE_CHAIN (c) = c2;
OMP_CLAUSE_DECL (c) = build_simple_mem_ref (t);
if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
OMP_CLAUSE_SIZE (c)
= TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (t)));
c = c2;
}
grp_start_p = pc;
grp_sentinel = OMP_CLAUSE_CHAIN (c);
tree nc = ai.expand_map_clause (c, OMP_CLAUSE_DECL (c),
addr_tokens, ort);
if (nc != error_mark_node)
c = nc;
}
}
break;
@ -9015,7 +8940,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
if (grp_start_p)
{
/* If we found a clause to remove, we want to remove the whole
expanded group, otherwise gimplify can get confused. */
expanded group, otherwise gimplify
(omp_resolve_clause_dependencies) can get confused. */
*grp_start_p = grp_sentinel;
pc = grp_start_p;
grp_start_p = NULL;
@ -10203,6 +10129,7 @@ finish_omp_target (location_t loc, tree clauses, tree body, bool combined_p)
case GOMP_MAP_ATTACH_DETACH:
case GOMP_MAP_ATTACH:
case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
case GOMP_MAP_POINTER:
case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
break;
default:

34
gcc/fortran/trans-openmp.cc
View File

@ -2405,8 +2405,8 @@ static vec<tree, va_heap, vl_embed> *doacross_steps;
static void
gfc_trans_omp_array_section (stmtblock_t *block, gfc_exec_op op,
gfc_omp_namelist *n, tree decl, bool element,
gomp_map_kind ptr_kind, tree &node, tree &node2,
tree &node3, tree &node4)
bool openmp, gomp_map_kind ptr_kind, tree &node,
tree &node2, tree &node3, tree &node4)
{
gfc_se se;
tree ptr, ptr2;
@ -2529,7 +2529,7 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_exec_op op,
struct - and adding an 'alloc: for the 'desc.data' pointer, which
would break as the 'desc' (the descriptor) is also mapped
(see node4 above). */
if (ptr_kind == GOMP_MAP_ATTACH_DETACH)
if (ptr_kind == GOMP_MAP_ATTACH_DETACH && !openmp)
STRIP_NOPS (OMP_CLAUSE_DECL (node3));
}
else
@ -2547,7 +2547,7 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_exec_op op,
decl, offset, NULL_TREE, NULL_TREE);
OMP_CLAUSE_DECL (node) = offset;
if (ptr_kind == GOMP_MAP_ALWAYS_POINTER)
if (ptr_kind == GOMP_MAP_ATTACH_DETACH && openmp)
return;
}
else
@ -3538,8 +3538,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
&& !(POINTER_TYPE_P (type)
&& GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (type))))
k = GOMP_MAP_FIRSTPRIVATE_POINTER;
gfc_trans_omp_array_section (block, op, n, decl, element, k,
node, node2, node3, node4);
gfc_trans_omp_array_section (block, op, n, decl, element,
!openacc, k, node, node2,
node3, node4);
}
else if (n->expr
&& n->expr->expr_type == EXPR_VARIABLE
@ -3578,10 +3579,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
node2 = build_omp_clause (input_location,
OMP_CLAUSE_MAP);
gomp_map_kind kind
= (openacc ? GOMP_MAP_ATTACH_DETACH
: GOMP_MAP_ALWAYS_POINTER);
OMP_CLAUSE_SET_MAP_KIND (node2, kind);
OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_ATTACH_DETACH);
OMP_CLAUSE_DECL (node2)
= POINTER_TYPE_P (TREE_TYPE (se.expr))
? se.expr
@ -3599,6 +3597,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
fold_convert (size_type_node,
se.string_length),
TYPE_SIZE_UNIT (tmp));
gomp_map_kind kind;
if (n->u.map_op == OMP_MAP_DELETE)
kind = GOMP_MAP_DELETE;
else if (op == EXEC_OMP_TARGET_EXIT_DATA)
@ -3685,9 +3684,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
node2 = build_omp_clause (input_location,
OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (node2,
openacc
? GOMP_MAP_ATTACH_DETACH
: GOMP_MAP_ALWAYS_POINTER);
GOMP_MAP_ATTACH_DETACH);
OMP_CLAUSE_DECL (node2) = build_fold_addr_expr (data);
OMP_CLAUSE_SIZE (node2) = size_int (0);
}
@ -3788,9 +3785,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
node3 = build_omp_clause (input_location,
OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (node3,
openacc
? GOMP_MAP_ATTACH_DETACH
: GOMP_MAP_ALWAYS_POINTER);
GOMP_MAP_ATTACH_DETACH);
OMP_CLAUSE_DECL (node3)
= gfc_conv_descriptor_data_get (inner);
/* Similar to gfc_trans_omp_array_section (details
@ -3813,11 +3808,10 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
{
/* An array element or section. */
bool element = lastref->u.ar.type == AR_ELEMENT;
gomp_map_kind kind = (openacc ? GOMP_MAP_ATTACH_DETACH
: GOMP_MAP_ALWAYS_POINTER);
gomp_map_kind kind = GOMP_MAP_ATTACH_DETACH;
gfc_trans_omp_array_section (block, op, n, inner, element,
kind, node, node2, node3,
node4);
!openacc, kind, node, node2,
node3, node4);
}
else
gcc_unreachable ();

1079
gcc/gimplify.cc
View File

File diff suppressed because it is too large Load Diff

425
gcc/omp-general.cc
View File

@ -45,6 +45,8 @@ along with GCC; see the file COPYING3. If not see
#include "data-streamer.h"
#include "streamer-hooks.h"
#include "opts.h"
#include "omp-general.h"
#include "tree-pretty-print.h"
enum omp_requires omp_requires_mask;
@ -3155,4 +3157,427 @@ omp_runtime_api_call (const_tree fndecl)
return omp_runtime_api_procname (IDENTIFIER_POINTER (declname));
}
namespace omp_addr_tokenizer {
/* We scan an expression by recursive descent, and build a vector of
"omp_addr_token *" pointers representing a "parsed" version of the
expression. The grammar we use is something like this:
expr0::
expr [section-access]
expr::
structured-expr access-method
| array-base access-method
structured-expr::
structure-base component-selector
arbitrary-expr::
(anything else)
structure-base::
DECL access-method
| structured-expr access-method
| arbitrary-expr access-method
array-base::
DECL
| arbitrary-expr
access-method::
DIRECT
| REF
| POINTER
| REF_TO_POINTER
| POINTER_OFFSET
| REF_TO_POINTER_OFFSET
| INDEXED_ARRAY
| INDEXED_REF_TO_ARRAY
| index-expr
index-expr::
INDEX_EXPR access-method
component-selector::
component-selector COMPONENT_REF
| component-selector ARRAY_REF
| COMPONENT_REF
This tokenized form is then used both in parsing, for OpenMP clause
expansion (for C and C++) and in gimplify.cc for sibling-list handling
(for C, C++ and Fortran). */
omp_addr_token::omp_addr_token (token_type t, tree e)
: type(t), expr(e)
{
}
omp_addr_token::omp_addr_token (access_method_kinds k, tree e)
: type(ACCESS_METHOD), expr(e)
{
u.access_kind = k;
}
omp_addr_token::omp_addr_token (token_type t, structure_base_kinds k, tree e)
: type(t), expr(e)
{
u.structure_base_kind = k;
}
static bool
omp_parse_component_selector (tree *expr0)
{
tree expr = *expr0;
tree last_component = NULL_TREE;
while (TREE_CODE (expr) == COMPONENT_REF
|| TREE_CODE (expr) == ARRAY_REF)
{
if (TREE_CODE (expr) == COMPONENT_REF)
last_component = expr;
expr = TREE_OPERAND (expr, 0);
if (TREE_CODE (TREE_TYPE (expr)) == REFERENCE_TYPE)
break;
}
if (!last_component)
return false;
*expr0 = last_component;
return true;
}
/* This handles references that have had convert_from_reference called on
them, and also those that haven't. */
static bool
omp_parse_ref (tree *expr0)
{
tree expr = *expr0;
if (TREE_CODE (TREE_TYPE (expr)) == REFERENCE_TYPE)
return true;
else if ((TREE_CODE (expr) == INDIRECT_REF
|| (TREE_CODE (expr) == MEM_REF
&& integer_zerop (TREE_OPERAND (expr, 1))))
&& TREE_CODE (TREE_TYPE (TREE_OPERAND (expr, 0))) == REFERENCE_TYPE)
{
*expr0 = TREE_OPERAND (expr, 0);
return true;
}
return false;
}
static bool
omp_parse_pointer (tree *expr0, bool *has_offset)
{
tree expr = *expr0;
*has_offset = false;
if ((TREE_CODE (expr) == INDIRECT_REF
|| (TREE_CODE (expr) == MEM_REF
&& integer_zerop (TREE_OPERAND (expr, 1))))
&& TREE_CODE (TREE_TYPE (TREE_OPERAND (expr, 0))) == POINTER_TYPE)
{
expr = TREE_OPERAND (expr, 0);
/* The Fortran FE sometimes emits a no-op cast here. */
STRIP_NOPS (expr);
while (1)
{
if (TREE_CODE (expr) == COMPOUND_EXPR)
{
expr = TREE_OPERAND (expr, 1);
STRIP_NOPS (expr);
}
else if (TREE_CODE (expr) == SAVE_EXPR)
expr = TREE_OPERAND (expr, 0);
else if (TREE_CODE (expr) == POINTER_PLUS_EXPR)
{
*has_offset = true;
expr = TREE_OPERAND (expr, 0);
}
else
break;
}
STRIP_NOPS (expr);
*expr0 = expr;
return true;
}
return false;
}
static bool
omp_parse_access_method (tree *expr0, enum access_method_kinds *kind)
{
tree expr = *expr0;
bool has_offset;
if (omp_parse_ref (&expr))
*kind = ACCESS_REF;
else if (omp_parse_pointer (&expr, &has_offset))
{
if (omp_parse_ref (&expr))
*kind = has_offset ? ACCESS_REF_TO_POINTER_OFFSET
: ACCESS_REF_TO_POINTER;
else
*kind = has_offset ? ACCESS_POINTER_OFFSET : ACCESS_POINTER;
}
else if (TREE_CODE (expr) == ARRAY_REF)
{
while (TREE_CODE (expr) == ARRAY_REF)
expr = TREE_OPERAND (expr, 0);
if (omp_parse_ref (&expr))
*kind = ACCESS_INDEXED_REF_TO_ARRAY;
else
*kind = ACCESS_INDEXED_ARRAY;
}
else
*kind = ACCESS_DIRECT;
STRIP_NOPS (expr);
*expr0 = expr;
return true;
}
static bool
omp_parse_access_methods (vec<omp_addr_token *> &addr_tokens, tree *expr0)
{
tree expr = *expr0;
enum access_method_kinds kind;
tree am_expr;
if (omp_parse_access_method (&expr, &kind))
am_expr = expr;
if (TREE_CODE (expr) == INDIRECT_REF
|| TREE_CODE (expr) == MEM_REF
|| TREE_CODE (expr) == ARRAY_REF)
omp_parse_access_methods (addr_tokens, &expr);
addr_tokens.safe_push (new omp_addr_token (kind, am_expr));
*expr0 = expr;
return true;
}
static bool omp_parse_structured_expr (vec<omp_addr_token *> &, tree *);
static bool
omp_parse_structure_base (vec<omp_addr_token *> &addr_tokens,
tree *expr0, structure_base_kinds *kind,
vec<omp_addr_token *> &base_access_tokens,
bool allow_structured = true)
{
tree expr = *expr0;
if (allow_structured)
omp_parse_access_methods (base_access_tokens, &expr);
if (DECL_P (expr))
{
*kind = BASE_DECL;
return true;
}
if (allow_structured && omp_parse_structured_expr (addr_tokens, &expr))
{
*kind = BASE_COMPONENT_EXPR;
*expr0 = expr;
return true;
}
*kind = BASE_ARBITRARY_EXPR;
*expr0 = expr;
return true;
}
static bool
omp_parse_structured_expr (vec<omp_addr_token *> &addr_tokens, tree *expr0)
{
tree expr = *expr0;
tree base_component = NULL_TREE;
structure_base_kinds struct_base_kind;
auto_vec<omp_addr_token *> base_access_tokens;
if (omp_parse_component_selector (&expr))
base_component = expr;
else
return false;
gcc_assert (TREE_CODE (expr) == COMPONENT_REF);
expr = TREE_OPERAND (expr, 0);
tree structure_base = expr;
if (!omp_parse_structure_base (addr_tokens, &expr, &struct_base_kind,
base_access_tokens))
return false;
addr_tokens.safe_push (new omp_addr_token (STRUCTURE_BASE, struct_base_kind,
structure_base));
addr_tokens.safe_splice (base_access_tokens);
addr_tokens.safe_push (new omp_addr_token (COMPONENT_SELECTOR,
base_component));
*expr0 = expr;
return true;
}
static bool
omp_parse_array_expr (vec<omp_addr_token *> &addr_tokens, tree *expr0)
{
tree expr = *expr0;
structure_base_kinds s_kind;
auto_vec<omp_addr_token *> base_access_tokens;
if (!omp_parse_structure_base (addr_tokens, &expr, &s_kind,
base_access_tokens, false))
return false;
addr_tokens.safe_push (new omp_addr_token (ARRAY_BASE, s_kind, expr));
addr_tokens.safe_splice (base_access_tokens);
*expr0 = expr;
return true;
}
/* Return TRUE if the ACCESS_METHOD token at index 'i' has a further
ACCESS_METHOD chained after it (e.g., if we're processing an expression
containing multiple pointer indirections). */
bool
omp_access_chain_p (vec<omp_addr_token *> &addr_tokens, unsigned i)
{
gcc_assert (addr_tokens[i]->type == ACCESS_METHOD);
return (i + 1 < addr_tokens.length ()
&& addr_tokens[i + 1]->type == ACCESS_METHOD);
}
/* Return the address of the object accessed by the ACCESS_METHOD token
at 'i': either of the next access method's expr, or of EXPR if we're at
the end of the list of tokens. */
tree
omp_accessed_addr (vec<omp_addr_token *> &addr_tokens, unsigned i, tree expr)
{
if (i + 1 < addr_tokens.length ())
return build_fold_addr_expr (addr_tokens[i + 1]->expr);
else
return build_fold_addr_expr (expr);
}
} /* namespace omp_addr_tokenizer. */
bool
omp_parse_expr (vec<omp_addr_token *> &addr_tokens, tree expr)
{
using namespace omp_addr_tokenizer;
auto_vec<omp_addr_token *> expr_access_tokens;
if (!omp_parse_access_methods (expr_access_tokens, &expr))
return false;
if (omp_parse_structured_expr (addr_tokens, &expr))
;
else if (omp_parse_array_expr (addr_tokens, &expr))
;
else
return false;
addr_tokens.safe_splice (expr_access_tokens);
return true;
}
DEBUG_FUNCTION void
debug_omp_tokenized_addr (vec<omp_addr_token *> &addr_tokens,
bool with_exprs)
{
using namespace omp_addr_tokenizer;
const char *sep = with_exprs ? " " : "";
for (auto e : addr_tokens)
{
const char *pfx = "";
fputs (sep, stderr);
switch (e->type)
{
case COMPONENT_SELECTOR:
fputs ("component_selector", stderr);
break;
case ACCESS_METHOD:
switch (e->u.access_kind)
{
case ACCESS_DIRECT:
fputs ("access_direct", stderr);
break;
case ACCESS_REF:
fputs ("access_ref", stderr);
break;
case ACCESS_POINTER:
fputs ("access_pointer", stderr);
break;
case ACCESS_POINTER_OFFSET:
fputs ("access_pointer_offset", stderr);
break;
case ACCESS_REF_TO_POINTER:
fputs ("access_ref_to_pointer", stderr);
break;
case ACCESS_REF_TO_POINTER_OFFSET:
fputs ("access_ref_to_pointer_offset", stderr);
break;
case ACCESS_INDEXED_ARRAY:
fputs ("access_indexed_array", stderr);
break;
case ACCESS_INDEXED_REF_TO_ARRAY:
fputs ("access_indexed_ref_to_array", stderr);
break;
}
break;
case ARRAY_BASE:
case STRUCTURE_BASE:
pfx = e->type == ARRAY_BASE ? "array_" : "struct_";
switch (e->u.structure_base_kind)
{
case BASE_DECL:
fprintf (stderr, "%sbase_decl", pfx);
break;
case BASE_COMPONENT_EXPR:
fputs ("base_component_expr", stderr);
break;
case BASE_ARBITRARY_EXPR:
fprintf (stderr, "%sbase_arbitrary_expr", pfx);
break;
}
break;
}
if (with_exprs)
{
fputs (" [", stderr);
print_generic_expr (stderr, e->expr);
fputc (']', stderr);
sep = ",\n ";
}
else
sep = " ";
}
fputs ("\n", stderr);
}
#include "gt-omp-general.h"

69
gcc/omp-general.h
View File

@ -153,4 +153,73 @@ get_openacc_privatization_dump_flags ()
extern tree omp_build_component_ref (tree obj, tree field);
namespace omp_addr_tokenizer {
/* These are the ways of accessing a variable that have special-case handling
in the middle end (gimplify, omp-lower, etc.). */
/* These are the kinds of access that an ACCESS_METHOD token can represent. */
enum access_method_kinds
{
ACCESS_DIRECT,
ACCESS_REF,
ACCESS_POINTER,
ACCESS_REF_TO_POINTER,
ACCESS_POINTER_OFFSET,
ACCESS_REF_TO_POINTER_OFFSET,
ACCESS_INDEXED_ARRAY,
ACCESS_INDEXED_REF_TO_ARRAY
};
/* These are the kinds that a STRUCTURE_BASE or ARRAY_BASE (except
BASE_COMPONENT_EXPR) can represent. */
enum structure_base_kinds
{
BASE_DECL,
BASE_COMPONENT_EXPR,
BASE_ARBITRARY_EXPR
};
/* The coarse type for an address token. These can have subtypes for
ARRAY_BASE or STRUCTURE_BASE (structure_base_kinds) or ACCESS_METHOD
(access_method_kinds). */
enum token_type
{
ARRAY_BASE,
STRUCTURE_BASE,
COMPONENT_SELECTOR,
ACCESS_METHOD
};
/* The struct that forms a single token of an address expression as parsed by
omp_parse_expr. These are typically held in a vec after parsing. */
struct omp_addr_token
{
enum token_type type;
tree expr;
union
{
access_method_kinds access_kind;
structure_base_kinds structure_base_kind;
} u;
omp_addr_token (token_type, tree);
omp_addr_token (access_method_kinds, tree);
omp_addr_token (token_type, structure_base_kinds, tree);
};
extern bool omp_access_chain_p (vec<omp_addr_token *> &, unsigned);
extern tree omp_accessed_addr (vec<omp_addr_token *> &, unsigned, tree);
}
typedef omp_addr_tokenizer::omp_addr_token omp_addr_token;
extern bool omp_parse_expr (vec<omp_addr_token *> &, tree);
#endif /* GCC_OMP_GENERAL_H */

7
gcc/omp-low.cc
View File

@ -1604,10 +1604,13 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
{
/* If this is an offloaded region, an attach operation should
only exist when the pointer variable is mapped in a prior
clause.
clause. An exception is if we have a reference (to pointer):
in that case we should have mapped "*decl" in a previous
mapping instead of "decl". Skip the assertion in that case.
If we had an error, we may not have attempted to sort clauses
properly, so avoid the test. */
if (is_gimple_omp_offloaded (ctx->stmt)
if (TREE_CODE (TREE_TYPE (decl)) != REFERENCE_TYPE
&& is_gimple_omp_offloaded (ctx->stmt)
&& !seen_error ())
gcc_assert
(maybe_lookup_decl (decl, ctx)

2
gcc/testsuite/c-c++-common/gomp/clauses-2.c
View File

@ -11,7 +11,7 @@ foo (int *p, int q, struct S t, int i, int j, int k, int l)
bar (p);
#pragma omp target firstprivate (p), map (p[0]) /* { dg-error "appears more than once in data clauses" } */
bar (p);
#pragma omp target map (p[0]) map (p) /* { dg-error "appears both in data and map clauses" } */
#pragma omp target map (p[0]) map (p)
bar (p);
#pragma omp target map (p) , map (p[0])
bar (p);

2
gcc/testsuite/c-c++-common/gomp/target-50.c
View File

@ -17,7 +17,7 @@ int main()
#pragma omp target map(tofrom: tmp->arr[0:10]) map(to: tmp->arr)
{ }
/* { dg-final { scan-tree-dump-times {map\(struct:\*tmp \[len: 1\]\) map\(to:tmp[._0-9]*->arr \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(attach:tmp[._0-9]*->arr \[bias: 0\]\)} 2 "gimple" { target { ! { nvptx*-*-* amdgcn*-*-* } } } } } */
/* { dg-final { scan-tree-dump-times {map\(struct:\*tmp \[len: 1\]\) map\(alloc:tmp[._0-9]*->arr \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(attach:tmp[._0-9]*->arr \[bias: 0\]\)} 2 "gimple" { target { ! { nvptx*-*-* amdgcn*-*-* } } } } } */
return 0;
}

3
gcc/testsuite/c-c++-common/gomp/target-enter-data-1.c
View File

@ -21,4 +21,5 @@ void func (struct foo *f, int n, int m)
#pragma omp target enter data map (to: f->bars[n].vectors[:f->bars[n].num_vectors])
}
/* { dg-final { scan-tree-dump-times "map\\(to:\\*_\[0-9\]+ \\\[len: _\[0-9\]+\\\]\\) map\\(attach:\[^-\]+->vectors \\\[bias: \[^\]\]+\\\]\\)" 3 "gimple" } } */
/* { dg-final { scan-tree-dump-times {map\(struct:\*f \[len: 1\]\) map\(alloc:[a-z0-9\._]+->vectors \[len: 0\]\) map\(to:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:[a-z0-9\._]+->vectors \[bias: [^\]]+\]\) map\(attach:\*_[0-9]+ \[bias: _[0-9]+\]\)} 1 "gimple" } } */
/* { dg-final { scan-tree-dump-times {map\(struct:\*\(f->bars \+ \(sizetype\) \(\([^\)]+\) n \* 16\)\) \[len: 1\]\) map\(alloc:[a-z0-9\._]+->vectors \[len: 0\]\) map\(to:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:[a-z0-9\._]+->vectors \[bias: [^\]]+\]\)} 2 "gimple" } } */

3
gcc/testsuite/c-c++-common/gomp/target-implicit-map-2.c
View File

@ -49,4 +49,5 @@ main (void)
/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(tofrom:a \[len: [0-9]+\]\[implicit\]\)} "gimple" } } */
/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(tofrom:a \[len: [0-9]+\]\[implicit\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(attach:a\.ptr \[bias: 0\]\)} "gimple" } } */
/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(struct:a \[len: 1\]\) map\(alloc:a\.ptr \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(attach:a\.ptr \[bias: 0\]\)} "gimple" } } */
/* { dg-final { scan-tree-dump-not {map\(struct:a \[len: 1\]\) map\(alloc:a\.ptr \[len: 0\]\)} "gimple" } } */

23
gcc/testsuite/g++.dg/gomp/static-component-1.C Normal file
View File

@ -0,0 +1,23 @@
/* { dg-do compile } */
/* Types with static members should be mappable. */
struct A {
static int x[10];
};
struct B {
A a;
};
int
main (int argc, char *argv[])
{
B *b = new B;
#pragma omp target map(b->a)
;
B bb;
#pragma omp target map(bb.a)
;
delete b;
}

2
gcc/testsuite/gcc.dg/gomp/target-3.c
View File

@ -13,4 +13,4 @@ void foo (struct S *s)
#pragma omp target enter data map (alloc: s->a, s->b)
}
/* { dg-final { scan-tree-dump-times "map\\(struct:\\*s \\\[len: 2\\\]\\) map\\(alloc:s->a \\\[len: \[0-9\]+\\\]\\) map\\(alloc:s->b \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */
/* { dg-final { scan-tree-dump-times "map\\(struct:\\*s \\\[len: 2\\\]\\) map\\(alloc:s\[\\._0-9\]+->a \\\[len: \[0-9\]+\\\]\\) map\\(alloc:s\[\\._0-9\]+->b \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */

2
gcc/testsuite/gfortran.dg/gomp/map-9.f90
View File

@ -2,7 +2,7 @@
! PR fortran/108545
! { dg-final { scan-tree-dump "#pragma omp target enter data map\\(struct:x \\\[len: 1\\\]\\) map\\(always,to:x.a \\\[len: \[0-9\]+\\\]\\) map\\(to:MEM <integer\\(kind=4\\)\\\[0:\\\]> \\\[\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\)_\[0-9\]+] \\\[len: _\[0-9\]+\\\]\\) map\\(always_pointer:x.a.data \\\[pointer assign, bias: 0\\\]\\)" "omplower" } }
! { dg-final { scan-tree-dump "#pragma omp target enter data map\\(struct:x \\\[len: 1\\\]\\) map\\(always,to:x\.a \\\[len: \[0-9\]+\\\]\\) map\\(to:MEM <integer\\(kind=4\\)\\\[0:\\\]> \\\[\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\)_\[0-9\]+] \\\[len: _\[0-9\]+\\\]\\) map\\(attach:x\.a\.data \\\[bias: 0\\\]\\)" "omplower" } }
program p
type t

4
gcc/tree.h
View File

@ -1827,6 +1827,10 @@ class auto_suppress_location_wrappers
NOTE: this is different than OMP_CLAUSE_MAP_IMPLICIT. */
#define OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P(NODE) \
(OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.deprecated_flag)
/* Nonzero for an attach/detach node whose decl was explicitly mapped on the
same directive. */
#define OMP_CLAUSE_ATTACHMENT_MAPPING_ERASED(NODE) \
TREE_STATIC (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP))
/* Flag that 'OMP_CLAUSE_DECL (NODE)' is to be made addressable during OMP
lowering. */

38
libgomp/target.c
View File

@ -703,7 +703,7 @@ gomp_map_pointer (struct target_mem_desc *tgt, struct goacc_asyncqueue *aq,
if (n == NULL)
{
if (allow_zero_length_array_sections)
cur_node.tgt_offset = 0;
cur_node.tgt_offset = cur_node.host_start;
else
{
gomp_mutex_unlock (&devicep->lock);
@ -742,7 +742,7 @@ gomp_map_fields_existing (struct target_mem_desc *tgt,
cur_node.host_start = (uintptr_t) hostaddrs[i];
cur_node.host_end = cur_node.host_start + sizes[i];
splay_tree_key n2 = splay_tree_lookup (mem_map, &cur_node);
splay_tree_key n2 = gomp_map_0len_lookup (mem_map, &cur_node);
kind = get_kind (short_mapkind, kinds, i);
implicit = get_implicit (short_mapkind, kinds, i);
if (n2
@ -839,8 +839,20 @@ gomp_attach_pointer (struct gomp_device_descr *devicep,
if ((void *) target == NULL)
{
gomp_mutex_unlock (&devicep->lock);
gomp_fatal ("attempt to attach null pointer");
/* As a special case, allow attaching NULL host pointers. This
allows e.g. unassociated Fortran pointers to be mapped
properly. */
data = 0;
gomp_debug (1,
"%s: attaching NULL host pointer, target %p "
"(struct base %p)\n", __FUNCTION__, (void *) devptr,
(void *) (n->tgt->tgt_start + n->tgt_offset));
gomp_copy_host2dev (devicep, aq, (void *) devptr, (void *) &data,
sizeof (void *), true, cbufp);
return;
}
s.host_start = target + bias;
@ -851,9 +863,8 @@ gomp_attach_pointer (struct gomp_device_descr *devicep,
{
if (allow_zero_length_array_sections)
/* When allowing attachment to zero-length array sections, we
allow attaching to NULL pointers when the target region is not
mapped. */
data = 0;
copy the host pointer when the target region is not mapped. */
data = target;
else
{
gomp_mutex_unlock (&devicep->lock);
@ -1097,7 +1108,8 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
tgt->list[i].key = NULL;
if (!aq
&& gomp_to_device_kind_p (get_kind (short_mapkind, kinds, i)
& typemask))
& typemask)
&& sizes[i] != 0)
gomp_coalesce_buf_add (&cbuf,
tgt_size - cur_node.host_end
+ (uintptr_t) hostaddrs[i],
@ -1464,7 +1476,17 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
+ sizes[last];
if (tgt->list[first].key != NULL)
continue;
if (sizes[last] == 0)
cur_node.host_end++;
n = splay_tree_lookup (mem_map, &cur_node);
if (sizes[last] == 0)
cur_node.host_end--;
if (n == NULL && cur_node.host_start == cur_node.host_end)
{
gomp_mutex_unlock (&devicep->lock);
gomp_fatal ("Struct pointer member not mapped (%p)",
(void*) hostaddrs[first]);
}
if (n == NULL)
{
size_t align = (size_t) 1 << (kind >> rshift);

275
libgomp/testsuite/libgomp.c++/baseptrs-3.C Normal file
View File

@ -0,0 +1,275 @@
#include <cstdlib>
#include <cstring>
#include <cassert>
struct sa0
{
int *ptr;
};
struct sb0
{
int arr[10];
};
struct sc0
{
sa0 a;
sb0 b;
sc0 (sa0 &my_a, sb0 &my_b) : a(my_a), b(my_b) {}
};
void
foo0 ()
{
sa0 my_a;
sb0 my_b;
my_a.ptr = (int *) malloc (sizeof (int) * 10);
sc0 my_c(my_a, my_b);
memset (my_c.a.ptr, 0, sizeof (int) * 10);
#pragma omp target map (my_c.a.ptr, my_c.a.ptr[:10])
{
for (int i = 0; i < 10; i++)
my_c.a.ptr[i] = i;
}
for (int i = 0; i < 10; i++)
assert (my_c.a.ptr[i] == i);
memset (my_c.b.arr, 0, sizeof (int) * 10);
#pragma omp target map (my_c.b.arr[:10])
{
for (int i = 0; i < 10; i++)
my_c.b.arr[i] = i;
}
for (int i = 0; i < 10; i++)
assert (my_c.b.arr[i] == i);
free (my_a.ptr);
}
struct sa
{
int *ptr;
};
struct sb
{
int arr[10];
};
struct sc
{
sa &a;
sb &b;
sc (sa &my_a, sb &my_b) : a(my_a), b(my_b) {}
};
void
foo ()
{
sa my_a;
sb my_b;
my_a.ptr = (int *) malloc (sizeof (int) * 10);
sc my_c(my_a, my_b);
memset (my_c.a.ptr, 0, sizeof (int) * 10);
#pragma omp target map (my_c.a.ptr, my_c.a.ptr[:10])
{
for (int i = 0; i < 10; i++)
my_c.a.ptr[i] = i;
}
for (int i = 0; i < 10; i++)
assert (my_c.a.ptr[i] == i);
memset (my_c.b.arr, 0, sizeof (int) * 10);
#pragma omp target map (my_c.b.arr[:10])
{
for (int i = 0; i < 10; i++)
my_c.b.arr[i] = i;
}
for (int i = 0; i < 10; i++)
assert (my_c.b.arr[i] == i);
free (my_a.ptr);
}
void
bar ()
{
sa my_a;
sb my_b;
my_a.ptr = (int *) malloc (sizeof (int) * 10);
sc my_c(my_a, my_b);
sc &my_cref = my_c;
memset (my_cref.a.ptr, 0, sizeof (int) * 10);
#pragma omp target map (my_cref.a.ptr, my_cref.a.ptr[:10])
{
for (int i = 0; i < 10; i++)
my_cref.a.ptr[i] = i;
}
for (int i = 0; i < 10; i++)
assert (my_cref.a.ptr[i] == i);
memset (my_cref.b.arr, 0, sizeof (int) * 10);
#pragma omp target map (my_cref.b.arr[:10])
{
for (int i = 0; i < 10; i++)
my_cref.b.arr[i] = i;
}
for (int i = 0; i < 10; i++)
assert (my_cref.b.arr[i] == i);
free (my_a.ptr);
}
struct scp0
{
sa *a;
sb *b;
scp0 (sa *my_a, sb *my_b) : a(my_a), b(my_b) {}
};
void
foop0 ()
{
sa *my_a = new sa;
sb *my_b = new sb;
my_a->ptr = new int[10];
scp0 *my_c = new scp0(my_a, my_b);
memset (my_c->a->ptr, 0, sizeof (int) * 10);
#pragma omp target map (my_c->a, my_c->a[:1], my_c->a->ptr, my_c->a->ptr[:10])
{
for (int i = 0; i < 10; i++)
my_c->a->ptr[i] = i;
}
for (int i = 0; i < 10; i++)
assert (my_c->a->ptr[i] == i);
memset (my_c->b->arr, 0, sizeof (int) * 10);
#pragma omp target map (my_c->b, my_c->b[:1], my_c->b->arr[:10])
{
for (int i = 0; i < 10; i++)
my_c->b->arr[i] = i;
}
for (int i = 0; i < 10; i++)
assert (my_c->b->arr[i] == i);
delete[] my_a->ptr;
delete my_a;
delete my_b;
}
struct scp
{
sa *&a;
sb *&b;
scp (sa *&my_a, sb *&my_b) : a(my_a), b(my_b) {}
};
void
foop ()
{
sa *my_a = new sa;
sb *my_b = new sb;
my_a->ptr = new int[10];
scp *my_c = new scp(my_a, my_b);
memset (my_c->a->ptr, 0, sizeof (int) * 10);
#pragma omp target map (my_c->a, my_c->a[:1], my_c->a->ptr, my_c->a->ptr[:10])
{
for (int i = 0; i < 10; i++)
my_c->a->ptr[i] = i;
}
for (int i = 0; i < 10; i++)
assert (my_c->a->ptr[i] == i);
memset (my_c->b->arr, 0, sizeof (int) * 10);
#pragma omp target map (my_c->b, my_c->b[:1], my_c->b->arr[:10])
{
for (int i = 0; i < 10; i++)
my_c->b->arr[i] = i;
}
for (int i = 0; i < 10; i++)
assert (my_c->b->arr[i] == i);
delete[] my_a->ptr;
delete my_a;
delete my_b;
}
void
barp ()
{
sa *my_a = new sa;
sb *my_b = new sb;
my_a->ptr = new int[10];
scp *my_c = new scp(my_a, my_b);
scp *&my_cref = my_c;
memset (my_cref->a->ptr, 0, sizeof (int) * 10);
#pragma omp target map (my_cref->a, my_cref->a[:1], my_cref->a->ptr, \
my_cref->a->ptr[:10])
{
for (int i = 0; i < 10; i++)
my_cref->a->ptr[i] = i;
}
for (int i = 0; i < 10; i++)
assert (my_cref->a->ptr[i] == i);
memset (my_cref->b->arr, 0, sizeof (int) * 10);
#pragma omp target map (my_cref->b, my_cref->b[:1], my_cref->b->arr[:10])
{
for (int i = 0; i < 10; i++)
my_cref->b->arr[i] = i;
}
for (int i = 0; i < 10; i++)
assert (my_cref->b->arr[i] == i);
delete my_a->ptr;
delete my_a;
delete my_b;
}
int main (int argc, char *argv[])
{
foo0 ();
foo ();
bar ();
foop0 ();
foop ();
barp ();
return 0;
}

3154
libgomp/testsuite/libgomp.c++/baseptrs-4.C Normal file
View File

File diff suppressed because it is too large Load Diff

62
libgomp/testsuite/libgomp.c++/baseptrs-5.C Normal file
View File

@ -0,0 +1,62 @@
// { dg-do run }
#include <cstring>
#include <cassert>
struct sa
{
int *ptr;
int *ptr2;
};
struct sb
{
int arr[10];
};
struct scp
{
sa *&a;
sb *&b;
scp (sa *&my_a, sb *&my_b) : a(my_a), b(my_b) {}
};
int
main ()
{
sa *my_a = new sa;
sb *my_b = new sb;
my_a->ptr = new int[10];
my_a->ptr2 = new int[10];
scp *my_c = new scp(my_a, my_b);
memset (my_c->a->ptr, 0, sizeof (int) * 10);
memset (my_c->a->ptr2, 0, sizeof (int) * 10);
#pragma omp target map (my_c->a, \
my_c->a->ptr, my_c->a->ptr[:10], \
my_c->a->ptr2, my_c->a->ptr2[:10])
{
for (int i = 0; i < 10; i++)
{
my_c->a->ptr[i] = i;
my_c->a->ptr2[i] = i * 2;
}
}
for (int i = 0; i < 10; i++)
{
assert (my_c->a->ptr[i] == i);
assert (my_c->a->ptr2[i] == i * 2);
}
delete[] my_a->ptr;
delete[] my_a->ptr2;
delete my_a;
delete my_b;
delete my_c;
return 0;
}

70
libgomp/testsuite/libgomp.c++/baseptrs-8.C Normal file
View File

@ -0,0 +1,70 @@
/* { dg-do run } */
#include <stdlib.h>
#include <assert.h>
#define N 1024
#define M 64
int main (void)
{
int *a_orig[N];
int *(&a)[N] = a_orig;
for (int i = 0; i < N; i++)
a[i] = (int *) calloc (M, sizeof (int));
/* 'target enter data'/'target exit data' with array of pointers. */
#pragma omp target enter data map(alloc: a[0:N])
for (int i = 0; i < N; i++)
{
#pragma omp target enter data map(to: a[i][0:M])
}
#pragma omp target map(alloc: a)
{
for (int i = 0; i < N; i++)
for (int j = 0; j < M; j++)
a[i][j] = i + j;
}
for (int i = 0; i < N; i++)
{
#pragma omp target exit data map(release: a[i]) map(from: a[i][0:M])
}
#pragma omp target exit data map(release: a, a[0:N])
/* 'target data' with array of pointers. */
#pragma omp target data map(alloc: a[0:N])
{
#pragma omp target data map(tofrom: a[5][0:M])
{
#pragma omp target map(alloc: a)
{
for (int i = 0; i < M; i++)
a[5][i]++;
}
}
}
/* 'target' with array of pointers. */
#pragma omp target data map(alloc: a[0:N])
{
#pragma omp target map(tofrom: a[7][0:M])
{
for (int i = 0; i < M; i++)
a[7][i] += 2;
}
}
for (int i = 0; i < N; i++)
for (int j = 0; j < M; j++)
assert (a[i][j] == i + j + (i == 5) + 2 * (i == 7));
for (int i = 0; i < N; i++)
free (a[i]);
return 0;
}

57
libgomp/testsuite/libgomp.c++/baseptrs-9.C Normal file
View File

@ -0,0 +1,57 @@
/* { dg-do run } */
#include <stdlib.h>
#include <assert.h>
int main (void)
{
int **a_orig,i,j,n;
int **&a = a_orig;
j = 3;
n = 12;
a = (int **) calloc (32, sizeof (int *));
for (int x = 0; x < 32; x++)
a[x] = (int *) calloc (32, sizeof (int));
for (int i = 2; i < 32; i++)
{
#pragma omp target enter data map(a, a[2:30])
#pragma omp target enter data map(a[i][j:n])
#pragma omp target map(alloc: a)
{
for (int x = j; x < j + n; x++)
a[i][x]++;
}
#pragma omp target exit data map(a[i][j:n])
#pragma omp target data map(a, a[i][j:n])
{
#pragma omp target map(alloc: a)
{
for (int x = j; x < j + n; x++)
a[i][x]++;
}
}
#pragma omp target exit data map(a, a[2:30])
#pragma omp target data map(a, a[2:30])
{
#pragma omp target map(a[i][j:n])
{
for (int x = j; x < j + n; x++)
a[i][x]++;
}
}
}
for (int i = 0; i < 32; i++)
for (int j = 0; j < 32; j++)
if (i >= 2 && j >= 3 && j < 15)
assert (a[i][j] == 3);
else
assert (a[i][j] == 0);
return 0;
}

59
libgomp/testsuite/libgomp.c++/class-array-1.C Normal file
View File

@ -0,0 +1,59 @@
/* { dg-do run } */
#include <cassert>
#define N 1024
class M {
int array[N];
public:
M ()
{
for (int i = 0; i < N; i++)
array[i] = 0;
}
void incr_with_this (int c)
{
#pragma omp target map(this->array[:N])
for (int i = 0; i < N; i++)
array[i] += c;
}
void incr_without_this (int c)
{
#pragma omp target map(array[:N])
for (int i = 0; i < N; i++)
array[i] += c;
}
void incr_implicit (int c)
{
#pragma omp target
for (int i = 0; i < N; i++)
array[i] += c;
}
void check (int c)
{
for (int i = 0; i < N; i++)
assert (array[i] == c);
}
};
int
main (int argc, char *argv[])
{
M m;
m.check (0);
m.incr_with_this (3);
m.check (3);
m.incr_without_this (5);
m.check (8);
m.incr_implicit (2);
m.check (10);
return 0;
}

80
libgomp/testsuite/libgomp.c++/ref-mapping-1.C Normal file
View File

@ -0,0 +1,80 @@
/* { dg-do run } */
#include <cassert>
void test_ref ()
{
int c_orig = 5;
int &c = c_orig;
#pragma omp target map(tofrom: c)
{
c++;
}
assert (c == 6);
}
void test_ref_to_ptr ()
{
int val = 5;
int *ptr_orig = &val;
int *&ptr_ref = ptr_orig;
#pragma omp target map(tofrom: ptr_ref[0])
{
(*ptr_ref)++;
}
assert (val == 6);
}
void test_ref_to_array ()
{
int arr[1];
int (&arr_ref)[1] = arr;
arr_ref[0] = 5;
#pragma omp target map(tofrom: arr_ref[0:1])
{
arr_ref[0]++;
}
assert (arr_ref[0] == 6);
#pragma omp target map(tofrom: arr_ref[0])
{
arr_ref[0]++;
}
assert (arr_ref[0] == 7);
}
void test_ref_to_ptr_array ()
{
int *arr[1];
int *(&arr_ref)[1] = arr;
int val = 5;
arr_ref[0] = &val;
#pragma omp target data map(alloc: arr_ref, arr_ref[0])
{
#pragma omp target map(tofrom: arr_ref[0][0:1])
{
arr_ref[0][0]++;
}
}
assert (arr_ref[0][0] == 6);
}
int main ()
{
test_ref ();
test_ref_to_ptr ();
test_ref_to_array ();
test_ref_to_ptr_array ();
return 0;
}

32
libgomp/testsuite/libgomp.c++/target-48.C Normal file
View File

@ -0,0 +1,32 @@
#include <cstring>
#include <cassert>
struct s {
int (&a)[10];
s(int (&a0)[10]) : a(a0) {}
};
int
main (int argc, char *argv[])
{
int la[10];
s v(la);
memset (la, 0, sizeof la);
#pragma omp target enter data map(to: v)
/* This mapping must use GOMP_MAP_ATTACH_DETACH not GOMP_MAP_ALWAYS_POINTER,
else the host reference v.a will be corrupted on copy-out. */
#pragma omp target map(v.a[0:10])
{
v.a[5]++;
}
#pragma omp target exit data map(from: v)
assert (v.a[5] == 1);
return 0;
}

37
libgomp/testsuite/libgomp.c++/target-49.C Normal file
View File

@ -0,0 +1,37 @@
#include <cstring>
#include <cassert>
struct s {
int (&a)[10];
s(int (&a0)[10]) : a(a0) {}
};
int
main (int argc, char *argv[])
{
int la[10];
s v_real(la);
s *v = &v_real;
memset (la, 0, sizeof la);
#pragma omp target enter data map(to: v)
/* Copying the whole v[0] here DOES NOT WORK yet because the reference 'a' is
not copied "as if" it was mapped explicitly as a member. FIXME. */
#pragma omp target enter data map(to: v[0])
#pragma omp target
{
v->a[5]++;
}
#pragma omp target exit data map(release: v[0])
#pragma omp target exit data map(from: v)
assert (v->a[5] == 1);
return 0;
}
// { dg-xfail-run-if "TODO" { offload_device_nonshared_as } }

34
libgomp/testsuite/libgomp.c++/target-exit-data-reftoptr-1.C Normal file
View File

@ -0,0 +1,34 @@
#include <cstring>
#include <cassert>
struct S {
int *&ptr;
S(int *&ptr_) : ptr(ptr_) { }
};
int main()
{
int *orig = new int[100];
memset (orig, 0, sizeof (int) * 100);
S svar(orig);
#pragma omp target enter data map(to: svar.ptr, svar.ptr[10:80])
#pragma omp target
{
for (int i = 10; i < 90; i++)
svar.ptr[i]++;
}
#pragma omp target exit data map(release: svar.ptr) map(from: svar.ptr[10:80])
for (int i = 0; i < 100; i++)
assert (i >= 10 && i < 90 && svar.ptr[i] == 1
|| svar.ptr[i] == 0);
delete orig;
return 0;
}

5
libgomp/testsuite/libgomp.c++/target-lambda-1.C
View File

@ -2,6 +2,7 @@
#include <cstdlib>
#include <cstring>
#include <cstdint>
template <typename L>
void
@ -22,9 +23,11 @@ struct S
auto fn = [=](void) -> bool
{
bool mapped;
uintptr_t hostptr = (uintptr_t) ptr;
uintptr_t hostiptr = (uintptr_t) iptr;
#pragma omp target map(from:mapped)
{
mapped = (ptr != NULL && iptr != NULL);
mapped = (ptr != (int*) hostptr && iptr != (int*) hostiptr);
if (mapped)
{
for (int i = 0; i < len; i++)

11
libgomp/testsuite/libgomp.c++/target-this-3.C
View File

@ -2,6 +2,7 @@
#include <stdio.h>
#include <string.h>
#include <stdint.h>
extern "C" void abort ();
struct S
@ -15,12 +16,13 @@ struct S
bool set_ptr (int n)
{
bool mapped;
uintptr_t hostptr = (uintptr_t) ptr;
#pragma omp target map(from:mapped)
{
if (ptr != NULL)
if (ptr != (int *) hostptr)
for (int i = 0; i < ptr_len; i++)
ptr[i] = n;
mapped = (ptr != NULL);
mapped = (ptr != (int *) hostptr);
}
return mapped;
}
@ -28,12 +30,13 @@ struct S
bool set_refptr (int n)
{
bool mapped;
uintptr_t hostrefptr = (uintptr_t) refptr;
#pragma omp target map(from:mapped)
{
if (refptr != NULL)
if (refptr != (int *) hostrefptr)
for (int i = 0; i < refptr_len; i++)
refptr[i] = n;
mapped = (refptr != NULL);
mapped = (refptr != (int *) hostrefptr);
}
return mapped;
}

11
libgomp/testsuite/libgomp.c++/target-this-4.C
View File

@ -4,6 +4,7 @@
#include <cstdlib>
#include <cstring>
#include <cstdint>
struct T
{
@ -18,12 +19,13 @@ struct T
auto fn = [=](void) -> bool
{
bool mapped;
uintptr_t hostptr = (uintptr_t) ptr;
#pragma omp target map(from:mapped)
{
if (ptr)
if (ptr != (int *) hostptr)
for (int i = 0; i < ptr_len; i++)
ptr[i] = n;
mapped = (ptr != NULL);
mapped = (ptr != (int *) hostptr);
}
return mapped;
};
@ -35,12 +37,13 @@ struct T
auto fn = [=](void) -> bool
{
bool mapped;
uintptr_t hostrefptr = (uintptr_t) refptr;
#pragma omp target map(from:mapped)
{
if (refptr)
if (refptr != (int *) hostrefptr)
for (int i = 0; i < refptr_len; i++)
refptr[i] = n;
mapped = (refptr != NULL);
mapped = (refptr != (int *) hostrefptr);
}
return mapped;
};

50
libgomp/testsuite/libgomp.c-c++-common/baseptrs-1.c Normal file
View File

@ -0,0 +1,50 @@
#include <stdlib.h>
#include <string.h>
#include <assert.h>
#include <stdio.h>
#define N 32
typedef struct {
int x2[10][N];
} x1type;
typedef struct {
x1type x1[10];
} p2type;
typedef struct {
p2type *p2;
} p1type;
typedef struct {
p1type *p1;
} x0type;
typedef struct {
x0type x0[10];
} p0type;
int main(int argc, char *argv[])
{
p0type *p0;
int k1 = 0, k2 = 0, k3 = 0, n = N;
p0 = (p0type *) malloc (sizeof *p0);
p0->x0[0].p1 = (p1type *) malloc (sizeof *p0->x0[0].p1);
p0->x0[0].p1->p2 = (p2type *) malloc (sizeof *p0->x0[0].p1->p2);
memset (p0->x0[0].p1->p2, 0, sizeof *p0->x0[0].p1->p2);
#pragma omp target map(tofrom: p0->x0[k1].p1->p2[k2].x1[k3].x2[4][0:n]) \
map(to: p0->x0[k1].p1, p0->x0[k1].p1->p2) \
map(to: p0->x0[k1].p1[0])
{
for (int i = 0; i < n; i++)
p0->x0[k1].p1->p2[k2].x1[k3].x2[4][i] = i;
}
for (int i = 0; i < n; i++)
assert (i == p0->x0[k1].p1->p2[k2].x1[k3].x2[4][i]);
return 0;
}

70
libgomp/testsuite/libgomp.c-c++-common/baseptrs-2.c Normal file
View File

@ -0,0 +1,70 @@
#include <stdlib.h>
#include <string.h>
#include <assert.h>
#define N 32
typedef struct {
int arr[N];
int *ptr;
} sc;
typedef struct {
sc *c;
} sb;
typedef struct {
sb *b;
sc *c;
} sa;
int main (int argc, char *argv[])
{
sa *p;
p = (sa *) malloc (sizeof *p);
p->b = (sb *) malloc (sizeof *p->b);
p->b->c = (sc *) malloc (sizeof *p->b->c);
p->c = (sc *) malloc (sizeof *p->c);
p->b->c->ptr = (int *) malloc (N * sizeof (int));
p->c->ptr = (int *) malloc (N * sizeof (int));
for (int i = 0; i < N; i++)
{
p->b->c->ptr[i] = 0;
p->c->ptr[i] = 0;
p->b->c->arr[i] = 0;
p->c->arr[i] = 0;
}
#pragma omp target map(to: p->b, p->b[0], p->c, p->c[0], p->b->c, p->b->c[0]) \
map(to: p->b->c->ptr, p->c->ptr) \
map(tofrom: p->b->c->ptr[:N], p->c->ptr[:N])
{
for (int i = 0; i < N; i++)
{
p->b->c->ptr[i] = i;
p->c->ptr[i] = i * 2;
}
}
#pragma omp target map(to: p->b, p->b[0], p->b->c, p->c) \
map(tofrom: p->c[0], p->b->c[0])
{
for (int i = 0; i < N; i++)
{
p->b->c->arr[i] = i * 3;
p->c->arr[i] = i * 4;
}
}
for (int i = 0; i < N; i++)
{
assert (p->b->c->ptr[i] == i);
assert (p->c->ptr[i] == i * 2);
assert (p->b->c->arr[i] == i * 3);
assert (p->c->arr[i] == i * 4);
}
return 0;
}

69
libgomp/testsuite/libgomp.c-c++-common/baseptrs-6.c Normal file
View File

@ -0,0 +1,69 @@
/* { dg-do run } */
#include <stdlib.h>
#include <assert.h>
#define N 1024
#define M 64
int main (void)
{
int *a[N];
for (int i = 0; i < N; i++)
a[i] = (int *) calloc (M, sizeof (int));
/* 'target enter data'/'target exit data' with array of pointers. */
#pragma omp target enter data map(alloc: a[0:N])
for (int i = 0; i < N; i++)
{
#pragma omp target enter data map(to: a[i][0:M])
}
#pragma omp target map(alloc: a)
{
for (int i = 0; i < N; i++)
for (int j = 0; j < M; j++)
a[i][j] = i + j;
}
for (int i = 0; i < N; i++)
{
#pragma omp target exit data map(release: a[i]) map(from: a[i][0:M])
}
#pragma omp target exit data map(release: a, a[0:N])
/* 'target data' with array of pointers. */
#pragma omp target data map(alloc: a[0:N])
{
#pragma omp target data map(tofrom: a[5][0:M])
{
#pragma omp target map(alloc: a)
{
for (int i = 0; i < M; i++)
a[5][i]++;
}
}
}
/* 'target' with array of pointers. */
#pragma omp target data map(alloc: a[0:N])
{
#pragma omp target map(tofrom: a[7][0:M])
{
for (int i = 0; i < M; i++)
a[7][i] += 2;
}
}
for (int i = 0; i < N; i++)
for (int j = 0; j < M; j++)
assert (a[i][j] == i + j + (i == 5) + 2 * (i == 7));
for (int i = 0; i < N; i++)
free (a[i]);
return 0;
}

56
libgomp/testsuite/libgomp.c-c++-common/baseptrs-7.c Normal file
View File

@ -0,0 +1,56 @@
/* { dg-do run } */
#include <stdlib.h>
#include <assert.h>
int main (void)
{
int **a,i,j,n;
j = 3;
n = 12;
a = (int **) calloc (32, sizeof (int *));
for (int x = 0; x < 32; x++)
a[x] = (int *) calloc (32, sizeof (int));
for (int i = 2; i < 32; i++)
{
#pragma omp target enter data map(a, a[2:30])
#pragma omp target enter data map(a[i][j:n])
#pragma omp target map(alloc: a)
{
for (int x = j; x < j + n; x++)
a[i][x]++;
}
#pragma omp target exit data map(a[i][j:n])
#pragma omp target data map(a, a[i][j:n])
{
#pragma omp target map(alloc: a)
{
for (int x = j; x < j + n; x++)
a[i][x]++;
}
}
#pragma omp target exit data map(a, a[2:30])
#pragma omp target data map(a, a[0:32])
{
#pragma omp target map(a[i][j:n])
{
for (int x = j; x < j + n; x++)
a[i][x]++;
}
}
}
for (int i = 0; i < 32; i++)
for (int j = 0; j < 32; j++)
if (i >= 2 && j >= 3 && j < 15)
assert (a[i][j] == 3);
else
assert (a[i][j] == 0);
return 0;
}

60
libgomp/testsuite/libgomp.c-c++-common/ptr-attach-2.c Normal file
View File

@ -0,0 +1,60 @@
#include <stdlib.h>
struct blk { int x, y; };
struct L
{
#define N 10
struct {
int num_blocks[N];
struct blk * blocks[N];
} m;
};
void foo (struct L *l)
{
for (int i = 0; i < N; i++)
{
l->m.blocks[i] = (struct blk *) malloc (sizeof (struct blk) * N);
l->m.num_blocks[i] = N;
}
#pragma omp target enter data map(to:l[:1])
for (int i = 0; i < N; i++)
{
#pragma omp target enter data map(to:l->m.blocks[i][:l->m.num_blocks[i]])
}
#pragma omp target
{
for (int i = 0; i < N; i++)
for (int j = 0; j < N; j++)
{
l->m.blocks[i][j].x = i + j;
l->m.blocks[i][j].y = i * j;
}
}
for (int i = 0; i < N; i++)
{
#pragma omp target exit data map(from:l->m.blocks[i][:l->m.num_blocks[i]])
}
#pragma omp target exit data map(from:l[:1])
for (int i = 0; i < N; i++)
for (int j = 0; j < N; j++)
{
if (l->m.blocks[i][j].x != i + j)
abort ();
if (l->m.blocks[i][j].y != i * j)
abort ();
}
}
int main (void)
{
struct L l;
foo (&l);
return 0;
}

2
libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-2.c
View File

@ -42,5 +42,7 @@ main (void)
#pragma omp target exit data map(from:a.ptr, a.ptr[:N])
free (a.ptr);
return 0;
}

50
libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-5.c Normal file
View File

@ -0,0 +1,50 @@
#include <stdlib.h>
#define N 10
struct S
{
int a, b;
int *ptr;
int c, d;
};
int
main (void)
{
struct S a;
a.ptr = (int *) malloc (sizeof (int) * N);
for (int i = 0; i < N; i++)
a.ptr[i] = 0;
#pragma omp target enter data map(to: a.ptr)
#pragma omp target enter data map(to: a.ptr[:N])
#pragma omp target
for (int i = 0; i < N; i++)
a.ptr[i] += 1;
#pragma omp target update from(a.ptr[:N])
for (int i = 0; i < N; i++)
if (a.ptr[i] != 1)
abort ();
#pragma omp target map(a.ptr[:N])
for (int i = 0; i < N; i++)
a.ptr[i] += 1;
#pragma omp target update from(a.ptr[:N])
for (int i = 0; i < N; i++)
if (a.ptr[i] != 2)
abort ();
#pragma omp target exit data map(release: a.ptr[:N])
#pragma omp target exit data map(release: a.ptr)
free (a.ptr);
return 0;
}

36
libgomp/testsuite/libgomp.c-c++-common/target-map-zlas-1.c Normal file
View File

@ -0,0 +1,36 @@
#include <stdlib.h>
#define N 10
struct S
{
int a, b;
int *ptr;
int c, d;
};
int
main (void)
{
struct S a;
a.ptr = (int *) malloc (sizeof (int) * N);
for (int i = 0; i < N; i++)
a.ptr[i] = 0;
#pragma omp target enter data map(to: a.ptr[:N])
#pragma omp target map(a, a.ptr[:0])
for (int i = 0; i < N; i++)
a.ptr[i] += 1;
#pragma omp target exit data map(from: a.ptr[:N])
for (int i = 0; i < N; i++)
if (a.ptr[i] != 1)
abort ();
free (a.ptr);
return 0;
}

3
libgomp/testsuite/libgomp.fortran/struct-elem-map-1.f90
View File

@ -409,3 +409,6 @@ contains
end subroutine eight
end program main
! Fixed by the "Fortran pointers and member mappings" patch
! { dg-xfail-run-if TODO { offload_device_nonshared_as } }

10
libgomp/testsuite/libgomp.fortran/target-enter-data-6.f90
View File

@ -3,6 +3,16 @@
! - arrays with array descriptors
! For those, the array descriptor / string length must be mapped with 'to:'
! This test fails without the following additional patches:
!
! "OpenMP: Pointers and member mappings":
! https://gcc.gnu.org/pipermail/gcc-patches/2023-August/627898.html
!
! "OpenMP/OpenACC: Reorganise OMP map clause handling in gimplify.cc":
! https://gcc.gnu.org/pipermail/gcc-patches/2023-August/627900.html
!
! { dg-xfail-run-if TODO { offload_device_nonshared_as } }
program main
implicit none