@@ -1253,6 +1253,61 @@ extern void c_omp_mark_declare_variant (location_t, tree, tree);
extern const char *c_omp_map_clause_name (tree, bool);
extern void c_omp_adjust_map_clauses (tree, bool);
+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 bool mappable_type (tree t);
+ virtual void emit_unmappable_type_notes (tree) { }
+
+ bool check_clause (tree);
+ tree get_root_term (bool);
+
+ tree get_address () { return orig; }
+ tree get_deref_origin ();
+ bool component_access_p ();
+
+ bool has_indirections_p ()
+ {
+ if (!root_term)
+ get_root_term (false);
+ return indirections;
+ }
+
+ bool indir_component_ref_p ()
+ {
+ return component_access_p () && has_indirections_p ();
+ }
+
+ bool map_supported_p ();
+
+ static tree get_origin (tree);
+ static tree peel_components (tree);
+ static tree maybe_peel_ref (tree);
+ static tree get_base_pointer (tree);
+ tree get_base_pointer () { return get_base_pointer (orig); }
+ static tree get_base_pointer_tgt (tree);
+ tree get_base_pointer_tgt () { return get_base_pointer_tgt (orig); }
+ static tree get_attachment_point (tree);
+ tree get_attachment_point () { return get_attachment_point (orig); }
+ bool maybe_zero_length_array_section (tree);
+};
+
enum c_omp_directive_kind {
C_OMP_DIR_STANDALONE,
C_OMP_DIR_CONSTRUCT,
@@ -3113,6 +3113,274 @@ c_omp_adjust_map_clauses (tree clauses, bool is_target)
}
}
+tree
+c_omp_address_inspector::get_deref_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. Look through that
+ access. */
+ t = maybe_peel_ref (t);
+
+ /* Find base pointer for POINTER_PLUS_EXPR, etc. */
+ t = get_origin (t);
+
+ return t;
+}
+
+bool
+c_omp_address_inspector::component_access_p ()
+{
+ tree t = maybe_peel_ref (orig);
+
+ t = get_origin (t);
+
+ return TREE_CODE (t) == COMPONENT_REF;
+}
+
+bool
+c_omp_address_inspector::check_clause (tree clause)
+{
+ tree t = get_deref_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 ()
+ && !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;
+}
+
+tree
+c_omp_address_inspector::get_root_term (bool checking)
+{
+ if (root_term && !checking)
+ return root_term;
+
+ tree t = get_deref_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)
+ {
+ 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;
+}
+
+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 = get_deref_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;
+}
+
+bool
+c_omp_address_inspector::mappable_type (tree t)
+{
+ return lang_hooks.types.omp_mappable_type (t);
+}
+
+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;
+}
+
+tree
+c_omp_address_inspector::peel_components (tree t)
+{
+ while (TREE_CODE (t) == COMPOUND_EXPR)
+ t = TREE_OPERAND (t, 1);
+
+ while (TREE_CODE (t) == COMPONENT_REF
+ && (TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF
+ || TREE_CODE (TREE_OPERAND (t, 0)) == ARRAY_REF))
+ {
+ t = TREE_OPERAND (t, 0);
+
+ while (TREE_CODE (t) == ARRAY_REF)
+ t = TREE_OPERAND (t, 0);
+ }
+
+ return t;
+}
+
+tree
+c_omp_address_inspector::maybe_peel_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;
+}
+
+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_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;
+ }
+}
+
+tree
+c_omp_address_inspector::get_base_pointer (tree expr)
+{
+ expr = maybe_peel_ref (expr);
+
+ if (TREE_CODE (TREE_TYPE (expr)) == POINTER_TYPE)
+ return expr;
+
+ if (TREE_CODE (TREE_TYPE (expr)) == REFERENCE_TYPE
+ && TREE_CODE (TREE_TYPE (TREE_TYPE (expr))) == POINTER_TYPE)
+ return expr;
+
+ expr = peel_components (expr);
+
+ if (TREE_CODE (expr) == COMPONENT_REF
+ && (DECL_P (TREE_OPERAND (expr, 0))
+ || TREE_CODE (TREE_OPERAND (expr, 0)) == INDIRECT_REF
+ || (TREE_CODE (TREE_OPERAND (expr, 0)) == MEM_REF
+ && integer_zerop (TREE_OPERAND (expr, 1)))))
+ expr = TREE_OPERAND (expr, 0);
+
+ if ((TREE_CODE (expr) == INDIRECT_REF || TREE_CODE (expr) == MEM_REF)
+ && TREE_CODE (TREE_TYPE (TREE_OPERAND (expr, 0))) == POINTER_TYPE)
+ return TREE_OPERAND (expr, 0);
+
+ return NULL_TREE;
+}
+
+tree
+c_omp_address_inspector::get_base_pointer_tgt (tree expr)
+{
+ expr = maybe_peel_ref (expr);
+
+ if (TREE_CODE (TREE_TYPE (expr)) == POINTER_TYPE)
+ return get_origin (expr);
+
+ if (TREE_CODE (TREE_TYPE (expr)) == REFERENCE_TYPE)
+ return expr;
+
+ return peel_components (expr);
+}
+
+tree
+c_omp_address_inspector::get_attachment_point (tree expr)
+{
+ tree baseptr = get_base_pointer (expr);
+
+ if (!baseptr)
+ return NULL_TREE;
+
+ return get_origin (baseptr);
+}
+
static const struct c_omp_directive omp_directives[] = {
/* Keep this alphabetically sorted by the first word. Non-null second/third
if any should precede null ones. */
@@ -13212,6 +13212,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
{
if (error_operand_p (t))
return error_mark_node;
+ c_omp_address_inspector t_insp (OMP_CLAUSE_LOCATION (c), t);
ret = t;
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_AFFINITY
&& OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND
@@ -13221,59 +13222,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 (TREE_CODE (t) == INDIRECT_REF)
- {
- 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
- || TREE_CODE (t) == INDIRECT_REF
- || 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 (!t_insp.check_clause (c))
+ return error_mark_node;
+ else if (t_insp.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 = t_insp.get_root_term (true);
+ else
+ t = t_insp.get_deref_origin ();
+ if (t == error_mark_node)
+ return error_mark_node;
if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
{
if (DECL_P (t))
@@ -13803,45 +13762,60 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
if (size)
size = c_fully_fold (size, false, NULL);
OMP_CLAUSE_SIZE (c) = size;
- 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;
+
+ c_omp_address_inspector t_insp (OMP_CLAUSE_LOCATION (c), t);
+
+ tree c2 = NULL_TREE;
+ tree baseptr = t_insp.get_base_pointer ();
+
+ /* DECL_P base pointers of struct accesses are special: e.g., a
+ FIRSTPRIVATE_POINTER may be created for them after a
+ GOMP_MAP_STRUCT node in gimplify.cc (if the base isn't mapped
+ TOFROM or similar otherwise). So, we don't need to create an
+ extra node here in that case. */
+ if (t_insp.component_access_p () && (!baseptr || DECL_P (baseptr)))
+ return false;
+
gcc_assert (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FORCE_DEVICEPTR);
- switch (OMP_CLAUSE_MAP_KIND (c))
+
+ if (t_insp.maybe_zero_length_array_section (c))
+ OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
+
+ c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
+ tree t_origin = t_insp.get_origin (t);
+ /* Use FIRSTPRIVATE_POINTER for plain or offset decls, otherwise
+ ATTACH_DETACH. */
+ if (DECL_P (t_origin))
{
- 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;
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
+ OMP_CLAUSE_DECL (c2) = t_origin;
}
- 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);
+ {
+ tree attach_pt = t_insp.get_attachment_point ();
+
+ /* Similarly to above, if we reach here and have a DECL_P
+ attachment point, we have e.g. an offset pointer as the base
+ of a struct access. We shouldn't create an ATTACH_DETACH
+ node for the DECL_P origin of that offset pointer. Instead,
+ the access will be handled by the struct sibling list
+ handling code in gimplify.cc. */
+ if (DECL_P (attach_pt))
+ return false;
+
+ if (!c_mark_addressable (t_origin))
+ return false;
+
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
+ OMP_CLAUSE_DECL (c2) = attach_pt;
+ }
OMP_CLAUSE_MAP_IMPLICIT (c2) = OMP_CLAUSE_MAP_IMPLICIT (c);
- if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
- && !c_mark_addressable (t))
- return false;
- OMP_CLAUSE_DECL (c2) = t;
+ tree ptr = t_insp.get_base_pointer_tgt (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,
@@ -14914,56 +14888,51 @@ 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 t_insp (OMP_CLAUSE_LOCATION (c), t);
+
+ if (t_insp.component_access_p ()
+ && !t_insp.get_base_pointer ())
{
- do
- {
- t = TREE_OPERAND (t, 0);
- if (TREE_CODE (t) == MEM_REF
- || TREE_CODE (t) == INDIRECT_REF)
- {
- 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);
+ tree rt = t_insp.get_root_term (false);
+
+ /* If the root term isn't a decl, these checks aren't
+ useful. */
+ if (!DECL_P (rt))
+ break;
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);
+ "clauses", rt);
else if (ort == C_ORT_ACC)
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));
}
}
}
@@ -14998,96 +14967,36 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
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 (TREE_CODE (t) == INDIRECT_REF
- || 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);
- }
- 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
- || TREE_CODE (TREE_OPERAND (t, 0)) == INDIRECT_REF
- || 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 (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 (!lang_hooks.types.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
- || TREE_CODE (t) == INDIRECT_REF
- || 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)
+ {
+ c_omp_address_inspector t_insp (OMP_CLAUSE_LOCATION (c), t);
+
+ if (!t_insp.check_clause (c))
+ t = error_mark_node;
+ else if (t_insp.component_access_p ()
+ && OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_)
+ t = t_insp.get_root_term (true);
+ else
+ t = t_insp.get_deref_origin ();
+
+ if (t == error_mark_node)
+ {
+ remove = true;
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;
- }
- }
+ }
+
+ indir_component_ref_p = t_insp.indir_component_ref_p ();
+
+ if (t_insp.component_access_p ()
+ && (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;
+ }
+ }
+
if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
{
error_at (OMP_CLAUSE_LOCATION (c),
@@ -5026,6 +5026,55 @@ 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;
+ }
+
+ bool mappable_type (tree t)
+ {
+ return cp_omp_mappable_type (t);
+ }
+
+ void emit_unmappable_type_notes (tree t)
+ {
+ cp_omp_emit_unmappable_type_notes (t);
+ }
+
+ static bool ref_p (tree t)
+ {
+ return (TYPE_REF_P (TREE_TYPE (t))
+ || REFERENCE_REF_P (t));
+ }
+
+ bool ref_p () { return ref_p (orig); }
+
+ static tree get_ref (tree t)
+ {
+ if (!ref_p (t))
+ return error_mark_node;
+ else if (TYPE_REF_P (TREE_TYPE (t)))
+ return t;
+ else if (REFERENCE_REF_P (t))
+ return TREE_OPERAND (t, 0);
+ return error_mark_node;
+ }
+
+ tree get_ref () { return get_ref (orig); }
+};
+
/* 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
@@ -5056,59 +5105,22 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
{
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 (TREE_CODE (t) == INDIRECT_REF)
- {
- 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
- || TREE_CODE (t) == INDIRECT_REF
- || 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 t_insp (OMP_CLAUSE_LOCATION (c), t);
+ tree t_refto = t_insp.maybe_peel_ref (t);
+
+ if (!t_insp.check_clause (c))
+ return error_mark_node;
+ else if (t_insp.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 = t_insp.get_root_term (true);
+ else
+ t = t_insp.get_deref_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)
@@ -5651,75 +5663,81 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
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_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;
+
+ cp_omp_address_inspector t_insp (OMP_CLAUSE_LOCATION (c), t);
+
+ if (t_insp.maybe_zero_length_array_section (c))
+ OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
+
tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
OMP_CLAUSE_MAP);
- if (TREE_CODE (t) == COMPONENT_REF)
+
+ tree baseptr = t_insp.get_base_pointer ();
+
+ /* DECL_P base pointers of struct accesses are special: e.g., a
+ FIRSTPRIVATE_POINTER may be created for them after a
+ GOMP_MAP_STRUCT node in gimplify.cc (if the base isn't mapped
+ TOFROM or similar otherwise). So, we don't need to create an
+ extra node here in that case. */
+ if (t_insp.component_access_p ()
+ && (!baseptr || DECL_P (baseptr))
+ && !t_insp.ref_p ())
+ return false;
+
+ tree t_origin = t_insp.get_origin (t);
+
+ if (DECL_P (t_origin))
{
- 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 (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;
- }
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
+ OMP_CLAUSE_DECL (c2) = t_origin;
}
- else if (REFERENCE_REF_P (t)
- && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
+ else if (!baseptr && !t_insp.ref_p ())
{
- 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);
+ /* If it's not a DECL_P or a component access, we've hit a
+ case we cannot handle. Bail out. */
+ sorry_at (OMP_CLAUSE_LOCATION (c),
+ "unsupported map expression %qE",
+ OMP_CLAUSE_DECL (c));
+ return true;
+ }
+ else if (t_insp.ref_p ()
+ && !t_insp.get_attachment_point ())
+ {
+ /* The array base itself is a reference. */
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
+ OMP_CLAUSE_DECL (c2) = t_insp.get_ref ();
}
else
- OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
+ {
+ tree attach_pt = t_insp.get_attachment_point ();
+
+ /* Similarly to above, if we reach here and have a DECL_P
+ attachment point, we have e.g. an offset pointer as the
+ base of a struct access. We shouldn't create an
+ ATTACH_DETACH node for the DECL_P origin of that offset
+ pointer. Instead, the access will be handled by the
+ struct sibling list handling code in gimplify.cc. */
+ if (DECL_P (attach_pt))
+ return false;
+
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
+
+ if (TYPE_REF_P (TREE_TYPE (attach_pt)))
+ attach_pt = convert_from_reference (attach_pt);
+ OMP_CLAUSE_DECL (c2) = attach_pt;
+ }
+
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;
+ tree ptr = t_insp.get_base_pointer_tgt (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);
@@ -5730,32 +5748,9 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
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))))
- {
- 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;
- }
}
}
+
return false;
}
@@ -7931,58 +7926,51 @@ 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)
+
+ cp_omp_address_inspector t_insp (OMP_CLAUSE_LOCATION (c), t);
+
+ if (t_insp.component_access_p ()
+ && !t_insp.get_base_pointer ())
{
- do
- {
- t = TREE_OPERAND (t, 0);
- if (REFERENCE_REF_P (t))
- t = TREE_OPERAND (t, 0);
- if (TREE_CODE (t) == MEM_REF
- || TREE_CODE (t) == INDIRECT_REF)
- {
- 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);
+ tree rt = t_insp.get_root_term (false);
+
+ /* If the root term isn't a decl, these checks aren't
+ useful. */
+ if (!DECL_P (rt))
+ break;
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);
+ " clauses", rt);
else if (ort == C_ORT_ACC)
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));
}
}
}
@@ -8017,104 +8005,41 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
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)
- {
- 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;
- }
- while (TREE_CODE (t) == INDIRECT_REF
- || 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);
- }
- 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)) == INDIRECT_REF
- || 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))
+
+ {
+ cp_omp_address_inspector t_insp (OMP_CLAUSE_LOCATION (c), t);
+
+ tree t_refto = c_omp_address_inspector::maybe_peel_ref (t);
+
+ if (!t_insp.check_clause (c))
+ t = error_mark_node;
+ else if (t_insp.component_access_p ()
+ && OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_)
+ t = t_insp.get_root_term (true);
+ else
+ t = t_insp.get_deref_origin ();
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH)
+ OMP_CLAUSE_DECL (c) = t_refto;
+ if (type_dependent_expression_p (t_refto))
+ break;
+ if (t == error_mark_node)
+ {
+ remove = true;
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 (!cp_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)]);
- cp_omp_emit_unmappable_type_notes (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
- || TREE_CODE (t) == INDIRECT_REF
- || 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)
+ }
+ indir_component_ref_p = t_insp.indir_component_ref_p ();
+ if (t_insp.component_access_p ()
+ && (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)
{
OMP_CLAUSE_DECL (c) = finish_non_static_data_member (t, NULL_TREE,
NULL_TREE);
@@ -8152,6 +8077,9 @@ 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
new file mode 100644
@@ -0,0 +1,21 @@
+/* { dg-do compile } */
+
+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) // { dg-error "'b->B::a' does not have a mappable type in 'map' clause" }
+ ;
+ B bb;
+#pragma omp target map(bb.a) // { dg-error "'bb\.B::a' does not have a mappable type in 'map' clause" }
+ ;
+ delete b;
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}