extern tree c_omp_check_context_selector (location_t, tree);
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);
/* Return next tree in the chain for chain_next walking of tree nodes. */
static inline tree
}
return omp_clause_code_name[OMP_CLAUSE_CODE (clause)];
}
+
+/* Used to merge map clause information in c_omp_adjust_map_clauses. */
+struct map_clause
+{
+ tree clause;
+ bool firstprivate_ptr_p;
+ bool decl_mapped;
+ bool omp_declare_target;
+ map_clause (void) : clause (NULL_TREE), firstprivate_ptr_p (false),
+ 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. */
+void
+c_omp_adjust_map_clauses (tree clauses, bool is_target)
+{
+ if (!is_target)
+ {
+ /* If this is not a target construct, just turn firstprivate pointers
+ into attach/detach, the runtime will check and do the rest. */
+
+ for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+ && DECL_P (OMP_CLAUSE_DECL (c))
+ && 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;
+ }
+
+ hash_map<tree, map_clause> maps;
+
+ for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && 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:
+ (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.
+
+ Here we iterate through all map clauses collecting these cases,
+ and merge them with a hash_map to process below. */
+
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+ && POINTER_TYPE_P (TREE_TYPE (OMP_CLAUSE_DECL (c))))
+ {
+ tree ptr = OMP_CLAUSE_DECL (c);
+ map_clause &mc = maps.get_or_insert (ptr);
+ if (mc.clause == NULL_TREE)
+ mc.clause = c;
+ mc.firstprivate_ptr_p = true;
+
+ if (is_global_var (ptr)
+ && lookup_attribute ("omp declare target",
+ DECL_ATTRIBUTES (ptr)))
+ mc.omp_declare_target = true;
+ }
+ else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALLOC
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FROM
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TOFROM
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TO
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_FROM
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TOFROM)
+ {
+ map_clause &mc = maps.get_or_insert (OMP_CLAUSE_DECL (c));
+ mc.decl_mapped = true;
+ }
+ }
+
+ for (hash_map<tree, map_clause>::iterator i = maps.begin ();
+ i != maps.end (); ++i)
+ {
+ map_clause &mc = (*i).second;
+
+ if (mc.firstprivate_ptr_p
+ && (mc.decl_mapped || mc.omp_declare_target))
+ {
+ OMP_CLAUSE_SET_MAP_KIND (mc.clause, GOMP_MAP_ATTACH_DETACH);
+ c_common_mark_addressable_vec (OMP_CLAUSE_DECL (mc.clause));
+ }
+ }
+}
tree clauses
= c_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK,
"#pragma omp target data");
+ c_omp_adjust_map_clauses (clauses, false);
int map_seen = 0;
for (tree *pc = &clauses; *pc;)
{
break;
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_ALWAYS_POINTER:
+ case GOMP_MAP_ATTACH_DETACH:
break;
default:
map_seen |= 1;
tree clauses
= c_parser_omp_all_clauses (parser, OMP_TARGET_ENTER_DATA_CLAUSE_MASK,
"#pragma omp target enter data");
+ c_omp_adjust_map_clauses (clauses, false);
int map_seen = 0;
for (tree *pc = &clauses; *pc;)
{
break;
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_ALWAYS_POINTER:
+ case GOMP_MAP_ATTACH_DETACH:
break;
default:
map_seen |= 1;
tree clauses
= c_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK,
"#pragma omp target exit data");
-
+ c_omp_adjust_map_clauses (clauses, false);
int map_seen = 0;
for (tree *pc = &clauses; *pc;)
{
break;
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_ALWAYS_POINTER:
+ case GOMP_MAP_ATTACH_DETACH:
break;
default:
map_seen |= 1;
OMP_TARGET_CLAUSES (stmt)
= c_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK,
"#pragma omp target");
+ c_omp_adjust_map_clauses (OMP_TARGET_CLAUSES (stmt), true);
+
pc = &OMP_TARGET_CLAUSES (stmt);
keep_next_level ();
block = c_begin_compound_stmt (true);
case GOMP_MAP_ALLOC:
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_ALWAYS_POINTER:
+ case GOMP_MAP_ATTACH_DETACH:
break;
default:
error_at (OMP_CLAUSE_LOCATION (*pc),
if (ort != C_ORT_OMP && ort != C_ORT_ACC)
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER);
else if (TREE_CODE (t) == COMPONENT_REF)
- {
- gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH
- : GOMP_MAP_ALWAYS_POINTER;
- OMP_CLAUSE_SET_MAP_KIND (c2, k);
- }
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
else
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
break;
if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
{
- if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
+ if (bitmap_bit_p (&map_field_head, DECL_UID (t))
+ || (ort == C_ORT_OMP
+ && bitmap_bit_p (&map_head, DECL_UID (t))))
break;
}
}
else
bitmap_set_bit (&generic_head, DECL_UID (t));
}
- else if (bitmap_bit_p (&map_head, DECL_UID (t)))
+ else if (bitmap_bit_p (&map_head, DECL_UID (t))
+ && (ort != C_ORT_OMP
+ || !bitmap_bit_p (&map_field_head, DECL_UID (t))))
{
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
error_at (OMP_CLAUSE_LOCATION (c),
remove = true;
}
else if (bitmap_bit_p (&generic_head, DECL_UID (t))
- || bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
+ && ort == C_ORT_ACC)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%qD appears more than once in data clauses", t);
+ remove = true;
+ }
+ else if (bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
{
if (ort == C_ORT_ACC)
error_at (OMP_CLAUSE_LOCATION (c),
tree clauses
= cp_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK,
"#pragma omp target data", pragma_tok);
+ c_omp_adjust_map_clauses (clauses, false);
int map_seen = 0;
for (tree *pc = &clauses; *pc;)
{
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
case GOMP_MAP_ALWAYS_POINTER:
+ case GOMP_MAP_ATTACH_DETACH:
break;
default:
map_seen |= 1;
tree clauses
= cp_parser_omp_all_clauses (parser, OMP_TARGET_ENTER_DATA_CLAUSE_MASK,
"#pragma omp target enter data", pragma_tok);
+ c_omp_adjust_map_clauses (clauses, false);
int map_seen = 0;
for (tree *pc = &clauses; *pc;)
{
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
case GOMP_MAP_ALWAYS_POINTER:
+ case GOMP_MAP_ATTACH_DETACH:
break;
default:
map_seen |= 1;
tree clauses
= cp_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK,
"#pragma omp target exit data", pragma_tok);
+ c_omp_adjust_map_clauses (clauses, false);
int map_seen = 0;
for (tree *pc = &clauses; *pc;)
{
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
case GOMP_MAP_ALWAYS_POINTER:
+ case GOMP_MAP_ATTACH_DETACH:
break;
default:
map_seen |= 1;
OMP_TARGET_CLAUSES (stmt)
= cp_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK,
"#pragma omp target", pragma_tok);
+ c_omp_adjust_map_clauses (OMP_TARGET_CLAUSES (stmt), true);
+
pc = &OMP_TARGET_CLAUSES (stmt);
keep_next_level (true);
OMP_TARGET_BODY (stmt) = cp_parser_omp_structured_block (parser, if_p);
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
case GOMP_MAP_ALWAYS_POINTER:
+ case GOMP_MAP_ATTACH_DETACH:
break;
default:
error_at (OMP_CLAUSE_LOCATION (*pc),
if ((ort & C_ORT_OMP_DECLARE_SIMD) != C_ORT_OMP && ort != C_ORT_ACC)
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER);
else if (TREE_CODE (t) == COMPONENT_REF)
- {
- gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH
- : GOMP_MAP_ALWAYS_POINTER;
- OMP_CLAUSE_SET_MAP_KIND (c2, k);
- }
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
else if (REFERENCE_REF_P (t)
&& TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
{
OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (c3, OMP_CLAUSE_MAP_KIND (c2));
OMP_CLAUSE_DECL (c3) = ptr;
- if (OMP_CLAUSE_MAP_KIND (c2) == GOMP_MAP_ALWAYS_POINTER)
- OMP_CLAUSE_DECL (c2) = build_simple_mem_ref (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;
t = TREE_OPERAND (t, 0);
OMP_CLAUSE_DECL (c) = t;
}
- if (ort == C_ORT_ACC
+ if ((ort == C_ORT_ACC || ort == C_ORT_OMP)
&& TREE_CODE (t) == COMPONENT_REF
&& TREE_CODE (TREE_OPERAND (t, 0)) == INDIRECT_REF)
t = TREE_OPERAND (TREE_OPERAND (t, 0), 0);
t = TREE_OPERAND (t, 0);
if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
{
- if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
+ if (bitmap_bit_p (&map_field_head, DECL_UID (t))
+ || (ort == C_ORT_OMP
+ && bitmap_bit_p (&map_head, DECL_UID (t))))
goto handle_map_references;
}
}
bitmap_set_bit (&generic_head, DECL_UID (t));
}
else if (bitmap_bit_p (&map_head, DECL_UID (t))
- && (ort != C_ORT_ACC
- || !bitmap_bit_p (&map_field_head, DECL_UID (t))))
+ && !bitmap_bit_p (&map_field_head, DECL_UID (t)))
{
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in motion clauses", t);
- if (ort == C_ORT_ACC)
+ else if (ort == C_ORT_ACC)
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in data clauses", t);
else
remove = true;
}
else if (bitmap_bit_p (&generic_head, DECL_UID (t))
- || bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
+ && ort == C_ORT_ACC)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%qD appears more than once in data clauses", t);
+ remove = true;
+ }
+ else if (bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
{
if (ort == C_ORT_ACC)
error_at (OMP_CLAUSE_LOCATION (c),
&& (OMP_CLAUSE_MAP_KIND (c)
!= GOMP_MAP_FIRSTPRIVATE_REFERENCE)
&& (OMP_CLAUSE_MAP_KIND (c)
- != GOMP_MAP_ALWAYS_POINTER))
+ != GOMP_MAP_ALWAYS_POINTER)
+ && (OMP_CLAUSE_MAP_KIND (c)
+ != GOMP_MAP_ATTACH_DETACH))
{
tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
OMP_CLAUSE_MAP);
if (TREE_CODE (t) == COMPONENT_REF)
- {
- gomp_map_kind k
- = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH
- : GOMP_MAP_ALWAYS_POINTER;
- OMP_CLAUSE_SET_MAP_KIND (c2, k);
- }
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
else
OMP_CLAUSE_SET_MAP_KIND (c2,
GOMP_MAP_FIRSTPRIVATE_REFERENCE);
return base;
}
+/* Returns true if EXPR is or contains (as a sub-component) BASE_PTR. */
+
+static bool
+is_or_contains_p (tree expr, tree base_ptr)
+{
+ while (expr != base_ptr)
+ if (TREE_CODE (base_ptr) == COMPONENT_REF)
+ base_ptr = TREE_OPERAND (base_ptr, 0);
+ else
+ break;
+ return expr == base_ptr;
+}
+
+/* Implement OpenMP 5.x map ordering rules for target directives. There are
+ several rules, and with some level of ambiguity, hopefully we can at least
+ collect the complexity here in one place. */
+
+static void
+omp_target_reorder_clauses (tree *list_p)
+{
+ /* Collect refs to alloc/release/delete maps. */
+ auto_vec<tree, 32> ard;
+ tree *cp = list_p;
+ while (*cp != NULL_TREE)
+ if (OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_ALLOC
+ || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_RELEASE
+ || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_DELETE))
+ {
+ /* Unlink cp and push to ard. */
+ tree c = *cp;
+ tree nc = OMP_CLAUSE_CHAIN (c);
+ *cp = nc;
+ ard.safe_push (c);
+
+ /* Any associated pointer type maps should also move along. */
+ while (*cp != NULL_TREE
+ && OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_FIRSTPRIVATE_REFERENCE
+ || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_FIRSTPRIVATE_POINTER
+ || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_ATTACH_DETACH
+ || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_POINTER
+ || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_ALWAYS_POINTER
+ || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_TO_PSET))
+ {
+ c = *cp;
+ nc = OMP_CLAUSE_CHAIN (c);
+ *cp = nc;
+ ard.safe_push (c);
+ }
+ }
+ else
+ cp = &OMP_CLAUSE_CHAIN (*cp);
+
+ /* Link alloc/release/delete maps to the end of list. */
+ for (unsigned int i = 0; i < ard.length (); i++)
+ {
+ *cp = ard[i];
+ cp = &OMP_CLAUSE_CHAIN (ard[i]);
+ }
+ *cp = NULL_TREE;
+
+ /* OpenMP 5.0 requires that pointer variables are mapped before
+ its use as a base-pointer. */
+ auto_vec<tree *, 32> atf;
+ for (tree *cp = list_p; *cp; cp = &OMP_CLAUSE_CHAIN (*cp))
+ if (OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP)
+ {
+ /* Collect alloc, to, from, to/from clause tree pointers. */
+ gomp_map_kind k = OMP_CLAUSE_MAP_KIND (*cp);
+ if (k == GOMP_MAP_ALLOC
+ || k == GOMP_MAP_TO
+ || k == GOMP_MAP_FROM
+ || k == GOMP_MAP_TOFROM
+ || k == GOMP_MAP_ALWAYS_TO
+ || k == GOMP_MAP_ALWAYS_FROM
+ || k == GOMP_MAP_ALWAYS_TOFROM)
+ atf.safe_push (cp);
+ }
+
+ for (unsigned int i = 0; i < atf.length (); i++)
+ if (atf[i])
+ {
+ tree *cp = atf[i];
+ tree decl = OMP_CLAUSE_DECL (*cp);
+ if (TREE_CODE (decl) == INDIRECT_REF || TREE_CODE (decl) == MEM_REF)
+ {
+ tree base_ptr = TREE_OPERAND (decl, 0);
+ STRIP_TYPE_NOPS (base_ptr);
+ for (unsigned int j = i + 1; j < atf.length (); j++)
+ {
+ tree *cp2 = atf[j];
+ tree decl2 = OMP_CLAUSE_DECL (*cp2);
+ if (is_or_contains_p (decl2, base_ptr))
+ {
+ /* Move *cp2 to before *cp. */
+ tree c = *cp2;
+ *cp2 = OMP_CLAUSE_CHAIN (c);
+ OMP_CLAUSE_CHAIN (c) = *cp;
+ *cp = c;
+ atf[j] = NULL;
+ }
+ }
+ }
+ }
+}
+
/* Scan the OMP clauses in *LIST_P, installing mappings into a new
and previous omp contexts. */
break;
}
+ if (code == OMP_TARGET
+ || code == OMP_TARGET_DATA
+ || code == OMP_TARGET_ENTER_DATA
+ || code == OMP_TARGET_EXIT_DATA)
+ omp_target_reorder_clauses (list_p);
+
while ((c = *list_p) != NULL)
{
bool remove = false;
}
else if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
|| (OMP_CLAUSE_MAP_KIND (c)
- == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
+ == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
&& TREE_CODE (OMP_CLAUSE_SIZE (c)) != INTEGER_CST)
{
OMP_CLAUSE_SIZE (c)
= get_initialized_tmp_var (OMP_CLAUSE_SIZE (c), pre_p, NULL,
false);
- omp_add_variable (ctx, OMP_CLAUSE_SIZE (c),
- GOVD_FIRSTPRIVATE | GOVD_SEEN);
+ if ((region_type & ORT_TARGET) != 0)
+ omp_add_variable (ctx, OMP_CLAUSE_SIZE (c),
+ GOVD_FIRSTPRIVATE | GOVD_SEEN);
}
+
if (!DECL_P (decl))
{
tree d = decl, *pd;
bool indir_p = false;
tree orig_decl = decl;
tree decl_ref = NULL_TREE;
- if ((region_type & ORT_ACC) != 0
+ if ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA)) != 0
&& TREE_CODE (*pd) == COMPONENT_REF
&& OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH
&& code != OACC_UPDATE)
while (TREE_CODE (decl) == COMPONENT_REF)
{
decl = TREE_OPERAND (decl, 0);
- if ((TREE_CODE (decl) == MEM_REF
- && integer_zerop (TREE_OPERAND (decl, 1)))
- || INDIRECT_REF_P (decl))
+ if (((TREE_CODE (decl) == MEM_REF
+ && integer_zerop (TREE_OPERAND (decl, 1)))
+ || INDIRECT_REF_P (decl))
+ && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
+ == POINTER_TYPE))
{
indir_p = true;
decl = TREE_OPERAND (decl, 0);
}
if (decl != orig_decl && DECL_P (decl) && indir_p)
{
- gomp_map_kind k = (code == OACC_EXIT_DATA) ? GOMP_MAP_DETACH
- : GOMP_MAP_ATTACH;
+ gomp_map_kind k
+ = ((code == OACC_EXIT_DATA || code == OMP_TARGET_EXIT_DATA)
+ ? GOMP_MAP_DETACH : GOMP_MAP_ATTACH);
/* We have a dereference of a struct member. Make this an
attach/detach operation, and ensure the base pointer is
mapped as a FIRSTPRIVATE_POINTER. */
tree next_clause = OMP_CLAUSE_CHAIN (c);
if (k == GOMP_MAP_ATTACH
&& code != OACC_ENTER_DATA
+ && code != OMP_TARGET_ENTER_DATA
&& (!next_clause
|| (OMP_CLAUSE_CODE (next_clause) != OMP_CLAUSE_MAP)
|| (OMP_CLAUSE_MAP_KIND (next_clause)
if (code == OACC_UPDATE
&& OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER);
- if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue)
- == GS_ERROR)
- {
- remove = true;
- break;
- }
if (DECL_P (decl)
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH
- && code != OACC_UPDATE)
+ && code != OACC_UPDATE
+ && code != OMP_TARGET_UPDATE)
{
if (error_operand_p (decl))
{
bool has_attachments = false;
/* For OpenACC, pointers in structs should trigger an
attach action. */
- if (attach_detach && (region_type & ORT_ACC) != 0)
+ if (attach_detach
+ && ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA))
+ || code == OMP_TARGET_ENTER_DATA
+ || code == OMP_TARGET_EXIT_DATA))
+
{
/* Turn a GOMP_MAP_ATTACH_DETACH clause into a
GOMP_MAP_ATTACH or GOMP_MAP_DETACH clause after we
have detected a case that needs a GOMP_MAP_STRUCT
mapping added. */
gomp_map_kind k
- = (code == OACC_EXIT_DATA) ? GOMP_MAP_DETACH
- : GOMP_MAP_ATTACH;
+ = ((code == OACC_EXIT_DATA || code == OMP_TARGET_EXIT_DATA)
+ ? GOMP_MAP_DETACH : GOMP_MAP_ATTACH);
OMP_CLAUSE_SET_MAP_KIND (c, k);
has_attachments = true;
}
break;
if (scp)
continue;
- tree d1 = OMP_CLAUSE_DECL (*sc);
- tree d2 = OMP_CLAUSE_DECL (c);
- while (TREE_CODE (d1) == ARRAY_REF)
- d1 = TREE_OPERAND (d1, 0);
- while (TREE_CODE (d2) == ARRAY_REF)
- d2 = TREE_OPERAND (d2, 0);
- if (TREE_CODE (d1) == INDIRECT_REF)
- d1 = TREE_OPERAND (d1, 0);
- if (TREE_CODE (d2) == INDIRECT_REF)
- d2 = TREE_OPERAND (d2, 0);
- while (TREE_CODE (d1) == COMPONENT_REF)
- if (TREE_CODE (d2) == COMPONENT_REF
- && TREE_OPERAND (d1, 1)
- == TREE_OPERAND (d2, 1))
- {
+ if ((region_type & ORT_ACC) != 0)
+ {
+ /* This duplicate checking code is currently only
+ enabled for OpenACC. */
+ tree d1 = OMP_CLAUSE_DECL (*sc);
+ tree d2 = OMP_CLAUSE_DECL (c);
+ while (TREE_CODE (d1) == ARRAY_REF)
d1 = TREE_OPERAND (d1, 0);
+ while (TREE_CODE (d2) == ARRAY_REF)
d2 = TREE_OPERAND (d2, 0);
- }
- else
- break;
- if (d1 == d2)
- {
- error_at (OMP_CLAUSE_LOCATION (c),
- "%qE appears more than once in map "
- "clauses", OMP_CLAUSE_DECL (c));
- remove = true;
- break;
+ if (TREE_CODE (d1) == INDIRECT_REF)
+ d1 = TREE_OPERAND (d1, 0);
+ if (TREE_CODE (d2) == INDIRECT_REF)
+ d2 = TREE_OPERAND (d2, 0);
+ while (TREE_CODE (d1) == COMPONENT_REF)
+ if (TREE_CODE (d2) == COMPONENT_REF
+ && TREE_OPERAND (d1, 1)
+ == TREE_OPERAND (d2, 1))
+ {
+ d1 = TREE_OPERAND (d1, 0);
+ d2 = TREE_OPERAND (d2, 0);
+ }
+ else
+ break;
+ if (d1 == d2)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%qE appears more than once in map "
+ "clauses", OMP_CLAUSE_DECL (c));
+ remove = true;
+ break;
+ }
}
if (maybe_lt (offset1, offsetn)
|| (known_eq (offset1, offsetn)
}
}
}
+
+ if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue)
+ == GS_ERROR)
+ {
+ remove = true;
+ break;
+ }
+
if (!remove
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH
break;
}
+ else
+ {
+ /* DECL_P (decl) == true */
+ tree *sc;
+ if (struct_map_to_clause
+ && (sc = struct_map_to_clause->get (decl)) != NULL
+ && OMP_CLAUSE_MAP_KIND (*sc) == GOMP_MAP_STRUCT
+ && decl == OMP_CLAUSE_DECL (*sc))
+ {
+ /* We have found a map of the whole structure after a
+ leading GOMP_MAP_STRUCT has been created, so refill the
+ leading clause into a map of the whole structure
+ variable, and remove the current one.
+ TODO: we should be able to remove some maps of the
+ following structure element maps if they are of
+ compatible TO/FROM/ALLOC type. */
+ OMP_CLAUSE_SET_MAP_KIND (*sc, OMP_CLAUSE_MAP_KIND (c));
+ OMP_CLAUSE_SIZE (*sc) = unshare_expr (OMP_CLAUSE_SIZE (c));
+ remove = true;
+ break;
+ }
+ }
flags = GOVD_MAP | GOVD_EXPLICIT;
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TO
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TOFROM)
flags |= GOVD_MAP_ALWAYS_TO;
+
+ if ((code == OMP_TARGET
+ || code == OMP_TARGET_DATA
+ || code == OMP_TARGET_ENTER_DATA
+ || code == OMP_TARGET_EXIT_DATA)
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
+ {
+ for (struct gimplify_omp_ctx *octx = outer_ctx; octx;
+ octx = octx->outer_context)
+ {
+ splay_tree_node n
+ = splay_tree_lookup (octx->variables,
+ (splay_tree_key) OMP_CLAUSE_DECL (c));
+ /* If this is contained in an outer OpenMP region as a
+ firstprivate value, remove the attach/detach. */
+ if (n && (n->value & GOVD_FIRSTPRIVATE))
+ {
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FIRSTPRIVATE_POINTER);
+ goto do_add;
+ }
+ }
+
+ enum gomp_map_kind map_kind = (code == OMP_TARGET_EXIT_DATA
+ ? GOMP_MAP_DETACH
+ : GOMP_MAP_ATTACH);
+ OMP_CLAUSE_SET_MAP_KIND (c, map_kind);
+ }
+
goto do_add;
case OMP_CLAUSE_DEPEND:
== GF_OMP_TARGET_KIND_OACC_KERNELS));
}
+/* Return true if STMT corresponds to an OpenMP target region. */
+static bool
+is_omp_target (gimple *stmt)
+{
+ if (gimple_code (stmt) == GIMPLE_OMP_TARGET)
+ {
+ int kind = gimple_omp_target_kind (stmt);
+ return (kind == GF_OMP_TARGET_KIND_REGION
+ || kind == GF_OMP_TARGET_KIND_DATA
+ || kind == GF_OMP_TARGET_KIND_ENTER_DATA
+ || kind == GF_OMP_TARGET_KIND_EXIT_DATA);
+ }
+ return false;
+}
+
/* If DECL is the artificial dummy VAR_DECL created for non-static
data member privatization, return the underlying "this" parameter,
otherwise return NULL. */
&& DECL_P (decl)
&& ((OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER
&& (OMP_CLAUSE_MAP_KIND (c)
- != GOMP_MAP_FIRSTPRIVATE_REFERENCE))
+ != GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+ && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH
+ && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH)
|| TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_TO
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_FROM
&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c))
break;
}
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && DECL_P (decl)
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
+ && is_omp_target (ctx->stmt))
+ {
+ /* If this is an offloaded region, an attach operation should
+ only exist when the pointer variable is mapped in a prior
+ clause. */
+ if (is_gimple_omp_offloaded (ctx->stmt))
+ gcc_assert
+ (maybe_lookup_decl (decl, ctx)
+ || (is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
+ && lookup_attribute ("omp declare target",
+ DECL_ATTRIBUTES (decl))));
+
+ /* By itself, attach/detach is generated as part of pointer
+ variable mapping and should not create new variables in the
+ offloaded region, however sender refs for it must be created
+ for its address to be passed to the runtime. */
+ tree field
+ = build_decl (OMP_CLAUSE_LOCATION (c),
+ FIELD_DECL, NULL_TREE, ptr_type_node);
+ SET_DECL_ALIGN (field, TYPE_ALIGN (ptr_type_node));
+ insert_field_into_struct (ctx->record_type, field);
+ /* To not clash with a map of the pointer variable itself,
+ attach/detach maps have their field looked up by the *clause*
+ tree expression, not the decl. */
+ gcc_assert (!splay_tree_lookup (ctx->field_map,
+ (splay_tree_key) c));
+ splay_tree_insert (ctx->field_map, (splay_tree_key) c,
+ (splay_tree_value) field);
+ break;
+ }
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
|| (OMP_CLAUSE_MAP_KIND (c)
&& is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
&& varpool_node::get_create (decl)->offloadable)
break;
+ if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
+ && is_omp_target (ctx->stmt)
+ && !is_gimple_omp_offloaded (ctx->stmt))
+ break;
if (DECL_P (decl))
{
if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
case GOMP_MAP_STRUCT:
case GOMP_MAP_ALWAYS_POINTER:
+ case GOMP_MAP_ATTACH:
+ case GOMP_MAP_DETACH:
break;
case GOMP_MAP_IF_PRESENT:
case GOMP_MAP_FORCE_ALLOC:
case GOMP_MAP_FORCE_DEVICEPTR:
case GOMP_MAP_DEVICE_RESIDENT:
case GOMP_MAP_LINK:
- case GOMP_MAP_ATTACH:
- case GOMP_MAP_DETACH:
case GOMP_MAP_FORCE_DETACH:
gcc_assert (is_gimple_omp_oacc (stmt));
break;
continue;
}
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
+ && is_omp_target (stmt))
+ {
+ gcc_assert (maybe_lookup_field (c, ctx));
+ map_cnt++;
+ continue;
+ }
+
if (!maybe_lookup_field (var, ctx))
continue;
gcc_assert (DECL_P (ovar2));
ovar = ovar2;
}
- if (!maybe_lookup_field (ovar, ctx))
+ if (!maybe_lookup_field (ovar, ctx)
+ && !(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)))
continue;
}
talign = TYPE_ALIGN_UNIT (TREE_TYPE (ovar));
if (DECL_P (ovar) && DECL_ALIGN_UNIT (ovar) > talign)
talign = DECL_ALIGN_UNIT (ovar);
- if (nc)
+
+ if (nc
+ && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
+ && is_omp_target (stmt))
+ {
+ var = lookup_decl_in_outer_ctx (ovar, ctx);
+ x = build_sender_ref (c, ctx);
+ gimplify_assign (x, build_fold_addr_expr (var), &ilist);
+ }
+ else if (nc)
{
var = lookup_decl_in_outer_ctx (ovar, ctx);
x = build_sender_ref (ovar, ctx);
bar (p);
#pragma omp target map (p[0]) map (p) /* { dg-error "appears both in data and map clauses" } */
bar (p);
- #pragma omp target map (p) , map (p[0]) /* { dg-error "appears both in data and map clauses" } */
+ #pragma omp target map (p) , map (p[0])
bar (p);
#pragma omp target map (q) map (q) /* { dg-error "appears more than once in map clauses" } */
bar (&q);
#pragma omp target map (p[0]) map (p[0]) /* { dg-error "appears more than once in data clauses" } */
bar (p);
- #pragma omp target map (t) map (t.r) /* { dg-error "appears more than once in map clauses" } */
+ #pragma omp target map (t) map (t.r)
bar (&t.r);
- #pragma omp target map (t.r) map (t) /* { dg-error "appears more than once in map clauses" } */
+ #pragma omp target map (t.r) map (t)
bar (&t.r);
- #pragma omp target map (t.r) map (t.r) /* { dg-error "appears more than once in map clauses" } */
+ #pragma omp target map (t.r) map (t.r)
bar (&t.r);
#pragma omp target firstprivate (t), map (t.r) /* { dg-error "appears both in data and map clauses" } */
bar (&t.r);
#pragma omp target map (t.r) firstprivate (t) /* { dg-error "appears both in data and map clauses" } */
bar (&t.r);
- #pragma omp target map (t.s[0]) map (t) /* { dg-error "appears more than once in map clauses" } */
+ #pragma omp target map (t.s[0]) map (t)
bar (t.s);
- #pragma omp target map (t) map(t.s[0]) /* { dg-error "appears more than once in map clauses" } */
+ #pragma omp target map (t) map(t.s[0])
bar (t.s);
#pragma omp target firstprivate (t) map (t.s[0]) /* { dg-error "appears both in data and map clauses" } */
bar (t.s);
#pragma omp target map (t.s[0]) firstprivate (t) /* { dg-error "appears both in data and map clauses" } */
bar (t.s);
- #pragma omp target map (t.s[0]) map (t.s[2]) /* { dg-error "appears more than once in map clauses" } */
+ #pragma omp target map (t.s[0]) map (t.s[2])
bar (t.s);
- #pragma omp target map (t.t[0:2]) map (t.t[4:6]) /* { dg-error "appears more than once in map clauses" } */
+ #pragma omp target map (t.t[0:2]) map (t.t[4:6])
bar (t.t);
- #pragma omp target map (t.t[i:j]) map (t.t[k:l]) /* { dg-error "appears more than once in map clauses" } */
+ #pragma omp target map (t.t[i:j]) map (t.t[k:l])
bar (t.t);
#pragma omp target map (t.s[0]) map (t.r)
bar (t.s);
#pragma omp target map (t.r) map (t) map (t.s[0]) firstprivate (t) /* { dg-error "appears both in data and map clauses" } */
bar (t.s);
#pragma omp target map (t) map (t.r) firstprivate (t) map (t.s[0]) /* { dg-error "appears both in data and map clauses" } */
- bar (t.s); /* { dg-error "appears more than once in map clauses" "" { target *-*-* } .-1 } */
+ bar (t.s);
}
--- /dev/null
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+void foo (void)
+{
+ /* Basic test to ensure to,from,tofrom is ordered before alloc,release,delete clauses. */
+ int a, b, c;
+ #pragma omp target enter data map(alloc:a) map(to:b) map(alloc:c)
+ #pragma omp target exit data map(from:a) map(release:b) map(from:c)
+
+ #pragma omp target map(alloc:a) map(tofrom:b) map(alloc:c)
+ a = b = c = 1;
+
+ #pragma omp target enter data map(to:a) map(alloc:b) map(to:c)
+ #pragma omp target exit data map(from:a) map(delete:b) map(from:c)
+}
+
+/* { dg-final { scan-tree-dump "pragma omp target enter data map\\(to:.* map\\(alloc:.* map\\(alloc:.*" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target exit data map\\(from:.* map\\(from:.* map\\(release:.*" "gimple" } } */
+
+/* { dg-final { scan-tree-dump "pragma omp target num_teams.* map\\(tofrom:.* map\\(alloc:.* map\\(alloc:.*" "gimple" } } */
+
+/* { dg-final { scan-tree-dump "pragma omp target enter data map\\(to:.* map\\(to:.* map\\(alloc:.*" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target exit data map\\(from:.* map\\(from:.* map\\(delete:.*" "gimple" } } */
integer :: i
end type t
type(t) v
-!$omp target enter data map(to:v%i, v%i) ! { dg-error "appears more than once in map clauses" }
+!$omp target enter data map(to:v%i, v%i)
end
/* Kind of the pragma, for which gomp_map_vars () is called. */
enum gomp_map_vars_kind
{
- GOMP_MAP_VARS_OPENACC,
- GOMP_MAP_VARS_TARGET,
- GOMP_MAP_VARS_DATA,
- GOMP_MAP_VARS_ENTER_DATA
+ GOMP_MAP_VARS_OPENACC = 1,
+ GOMP_MAP_VARS_TARGET = 2,
+ GOMP_MAP_VARS_DATA = 4,
+ GOMP_MAP_VARS_ENTER_DATA = 8
};
extern void gomp_acc_declare_allocate (bool, size_t, void **, size_t *,
struct target_mem_desc *tgt
= gomp_map_vars (acc_dev, mapnum, &hostaddrs, &devaddrs, &sizes,
- &kinds, true, GOMP_MAP_VARS_ENTER_DATA);
+ &kinds, true,
+ GOMP_MAP_VARS_OPENACC | GOMP_MAP_VARS_ENTER_DATA);
assert (tgt);
assert (tgt->list_count == 1);
splay_tree_key n = tgt->list[0].key;
struct target_mem_desc *tgt
= gomp_map_vars_async (acc_dev, aq, mapnum, hostaddrs, NULL, sizes,
- kinds, true, GOMP_MAP_VARS_ENTER_DATA);
+ kinds, true, (GOMP_MAP_VARS_OPENACC
+ | GOMP_MAP_VARS_ENTER_DATA));
assert (tgt);
assert (tgt->list_count == 1);
n = tgt->list[0].key;
struct target_mem_desc *tgt
= gomp_map_vars_async (acc_dev, aq, groupnum, &hostaddrs[i], NULL,
&sizes[i], &kinds[i], true,
- GOMP_MAP_VARS_ENTER_DATA);
+ (GOMP_MAP_VARS_OPENACC
+ | GOMP_MAP_VARS_ENTER_DATA));
assert (tgt);
gomp_mutex_lock (&acc_dev->lock);
struct target_mem_desc *tgt
= gomp_malloc (sizeof (*tgt) + sizeof (tgt->list[0]) * mapnum);
tgt->list_count = mapnum;
- tgt->refcount = pragma_kind == GOMP_MAP_VARS_ENTER_DATA ? 0 : 1;
+ tgt->refcount = (pragma_kind & GOMP_MAP_VARS_ENTER_DATA) ? 0 : 1;
tgt->device_descr = devicep;
tgt->prev = NULL;
struct gomp_coalesce_buf cbuf, *cbufp = NULL;
/* OpenACC 'attach'/'detach' doesn't affect
structured/dynamic reference counts ('n->refcount',
'n->dynamic_refcount'). */
+
+ gomp_attach_pointer (devicep, aq, mem_map, n,
+ (uintptr_t) hostaddrs[i], sizes[i],
+ cbufp);
}
- else
+ else if ((pragma_kind & GOMP_MAP_VARS_OPENACC) != 0)
{
gomp_mutex_unlock (&devicep->lock);
gomp_fatal ("outer struct not mapped for attach");
}
- gomp_attach_pointer (devicep, aq, mem_map, n,
- (uintptr_t) hostaddrs[i], sizes[i],
- cbufp);
continue;
}
default:
/* If the variable from "omp target enter data" map-list was already mapped,
tgt is not needed. Otherwise tgt will be freed by gomp_unmap_vars or
gomp_exit_data. */
- if (pragma_kind == GOMP_MAP_VARS_ENTER_DATA && tgt->refcount == 0)
+ if ((pragma_kind & GOMP_MAP_VARS_ENTER_DATA) && tgt->refcount == 0)
{
free (tgt);
tgt = NULL;
return;
}
+ for (i = 0; i < mapnum; i++)
+ if ((kinds[i] & typemask) == GOMP_MAP_DETACH)
+ {
+ struct splay_tree_key_s cur_node;
+ cur_node.host_start = (uintptr_t) hostaddrs[i];
+ cur_node.host_end = cur_node.host_start + sizeof (void *);
+ splay_tree_key n = splay_tree_lookup (&devicep->mem_map, &cur_node);
+
+ if (n)
+ gomp_detach_pointer (devicep, NULL, n, (uintptr_t) hostaddrs[i],
+ false, NULL);
+ }
+
for (i = 0; i < mapnum; i++)
{
struct splay_tree_key_s cur_node;
cur_node.host_end - cur_node.host_start);
if (k->refcount == 0)
gomp_remove_var (devicep, k);
+ break;
+ case GOMP_MAP_DETACH:
break;
default:
gomp_mutex_unlock (&devicep->lock);
&kinds[i], true, GOMP_MAP_VARS_ENTER_DATA);
i += j - i - 1;
}
+ else if (i + 1 < mapnum && (kinds[i + 1] & 0xff) == GOMP_MAP_ATTACH)
+ {
+ /* An attach operation must be processed together with the mapped
+ base-pointer list item. */
+ gomp_map_vars (devicep, 2, &hostaddrs[i], NULL, &sizes[i], &kinds[i],
+ true, GOMP_MAP_VARS_ENTER_DATA);
+ i += 1;
+ }
else
gomp_map_vars (devicep, 1, &hostaddrs[i], NULL, &sizes[i], &kinds[i],
true, GOMP_MAP_VARS_ENTER_DATA);
--- /dev/null
+#include <stdlib.h>
+
+struct S
+{
+ int a, b;
+ int *ptr;
+ int c, d;
+};
+typedef struct S S;
+
+#pragma omp declare target
+int *gp;
+#pragma omp end declare target
+
+#define N 10
+int main (void)
+{
+ /* Test to see if pointer attachment works, for scalar pointers,
+ and pointer fields in structures. */
+
+ int *ptr = (int *) malloc (sizeof (int) * N);
+ int *orig_ptr = ptr;
+
+ #pragma omp target map (ptr, ptr[:N])
+ {
+ for (int i = 0; i < N; i++)
+ ptr[i] = N - i;
+ }
+
+ if (ptr != orig_ptr)
+ abort ();
+
+ for (int i = 0; i < N; i++)
+ if (ptr[i] != N - i)
+ abort ();
+
+ S s = { 0 };
+ s.ptr = ptr;
+ #pragma omp target map (s, s.ptr[:N])
+ {
+ for (int i = 0; i < N; i++)
+ s.ptr[i] = i;
+
+ s.a = 1;
+ s.b = 2;
+ }
+
+ if (s.ptr != ptr)
+ abort ();
+
+ for (int i = 0; i < N; i++)
+ if (s.ptr[i] != i)
+ abort ();
+
+ if (s.a != 1 || s.b != 2 || s.c != 0 || s.d != 0)
+ abort ();
+
+ gp = (int *) malloc (sizeof (int) * N);
+ orig_ptr = gp;
+
+ for (int i = 0; i < N; i++)
+ gp[i] = i - 1;
+
+ #pragma omp target map (gp[:N])
+ {
+ for (int i = 0; i < N; i++)
+ gp[i] += 1;
+ }
+
+ if (gp != orig_ptr)
+ abort ();
+
+ for (int i = 0; i < N; i++)
+ if (gp[i] != i)
+ abort ();
+
+ free (ptr);
+ free (gp);
+
+ return 0;
+}
+
projects / gcc.git / commitdiff