case OMP_CLAUSE_DEVICE:
case OMP_CLAUSE_MAP:
case OMP_CLAUSE_IS_DEVICE_PTR:
+ case OMP_CLAUSE_HAS_DEVICE_ADDR:
case OMP_CLAUSE_DEFAULTMAP:
case OMP_CLAUSE_DEPEND:
s = C_OMP_CLAUSE_SPLIT_TARGET;
};
-/* All clauses defined by OpenACC 2.0, and OpenMP 2.5, 3.0, 3.1, 4.0, 4.5
- and 5.0. Used internally by both C and C++ parsers. */
+/* All clauses defined by OpenACC 2.0, and OpenMP 2.5, 3.0, 3.1, 4.0, 4.5, 5.0,
+ and 5.1. Used internally by both C and C++ parsers. */
enum pragma_omp_clause {
PRAGMA_OMP_CLAUSE_NONE = 0,
PRAGMA_OMP_CLAUSE_FOR,
PRAGMA_OMP_CLAUSE_FROM,
PRAGMA_OMP_CLAUSE_GRAINSIZE,
+ PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR,
PRAGMA_OMP_CLAUSE_HINT,
PRAGMA_OMP_CLAUSE_IF,
PRAGMA_OMP_CLAUSE_IN_REDUCTION,
result = PRAGMA_OMP_CLAUSE_GRAINSIZE;
break;
case 'h':
- if (!strcmp ("hint", p))
+ if (!strcmp ("has_device_addr", p))
+ result = PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR;
+ else if (!strcmp ("hint", p))
result = PRAGMA_OMP_CLAUSE_HINT;
else if (!strcmp ("host", p))
result = PRAGMA_OACC_CLAUSE_HOST;
case OMP_CLAUSE_REDUCTION:
case OMP_CLAUSE_IN_REDUCTION:
case OMP_CLAUSE_TASK_REDUCTION:
+ case OMP_CLAUSE_HAS_DEVICE_ADDR:
array_section_p = false;
dims.truncate (0);
while (c_parser_next_token_is (parser, CPP_OPEN_SQUARE))
list);
}
+/* OpenMP 5.1:
+ has_device_addr ( variable-list ) */
+
+static tree
+c_parser_omp_clause_has_device_addr (c_parser *parser, tree list)
+{
+ return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_HAS_DEVICE_ADDR,
+ list);
+}
+
/* OpenMP 4.5:
is_device_ptr ( variable-list ) */
clauses = c_parser_omp_clause_use_device_addr (parser, clauses);
c_name = "use_device_addr";
break;
+ case PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR:
+ clauses = c_parser_omp_clause_has_device_addr (parser, clauses);
+ c_name = "has_device_addr";
+ break;
case PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR:
clauses = c_parser_omp_clause_is_device_ptr (parser, clauses);
c_name = "is_device_ptr";
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEFAULTMAP) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IN_REDUCTION) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_THREAD_LIMIT) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR))
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR)\
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR))
static bool
c_parser_omp_target (c_parser *parser, enum pragma_context context, bool *if_p)
}
first = c_fully_fold (first, false, NULL);
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;
{
bitmap_head generic_head, firstprivate_head, lastprivate_head;
bitmap_head aligned_head, map_head, map_field_head, map_firstprivate_head;
- bitmap_head oacc_reduction_head;
+ bitmap_head oacc_reduction_head, is_on_device_head;
tree c, t, type, *pc;
tree simdlen = NULL_TREE, safelen = NULL_TREE;
bool branch_seen = false;
/* If ort == C_ORT_OMP used as nontemporal_head or use_device_xxx_head
instead and for ort == C_ORT_OMP_TARGET used as in_reduction_head. */
bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack);
+ bitmap_initialize (&is_on_device_head, &bitmap_default_obstack);
if (ort & C_ORT_ACC)
for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
"%qE appears more than once in data clauses", t);
remove = true;
}
- else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_PRIVATE
+ else if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_PRIVATE
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR)
&& bitmap_bit_p (&map_head, DECL_UID (t)))
{
if (ort == C_ORT_ACC)
"%qD appears more than once in data clauses", t);
remove = true;
}
- else if (bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
+ 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)
error_at (OMP_CLAUSE_LOCATION (c),
case OMP_CLAUSE_IS_DEVICE_PTR:
case OMP_CLAUSE_USE_DEVICE_PTR:
t = OMP_CLAUSE_DECL (c);
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR)
+ bitmap_set_bit (&is_on_device_head, DECL_UID (t));
if (TREE_CODE (TREE_TYPE (t)) != POINTER_TYPE)
{
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR
}
goto check_dup_generic;
+ case OMP_CLAUSE_HAS_DEVICE_ADDR:
+ t = OMP_CLAUSE_DECL (c);
+ if (TREE_CODE (t) == TREE_LIST)
+ {
+ if (handle_omp_array_sections (c, ort))
+ remove = true;
+ else
+ {
+ t = OMP_CLAUSE_DECL (c);
+ while (TREE_CODE (t) == ARRAY_REF)
+ t = TREE_OPERAND (t, 0);
+ }
+ }
+ bitmap_set_bit (&is_on_device_head, DECL_UID (t));
+ if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
+ c_mark_addressable (t);
+ goto check_dup_generic_t;
+
case OMP_CLAUSE_USE_DEVICE_ADDR:
t = OMP_CLAUSE_DECL (c);
if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
result = PRAGMA_OMP_CLAUSE_GRAINSIZE;
break;
case 'h':
- if (!strcmp ("hint", p))
+ if (!strcmp ("has_device_addr", p))
+ result = PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR;
+ else if (!strcmp ("hint", p))
result = PRAGMA_OMP_CLAUSE_HINT;
else if (!strcmp ("host", p))
result = PRAGMA_OACC_CLAUSE_HOST;
case OMP_CLAUSE_REDUCTION:
case OMP_CLAUSE_IN_REDUCTION:
case OMP_CLAUSE_TASK_REDUCTION:
+ case OMP_CLAUSE_HAS_DEVICE_ADDR:
array_section_p = false;
dims.truncate (0);
while (cp_lexer_next_token_is (parser->lexer, CPP_OPEN_SQUARE))
clauses);
c_name = "is_device_ptr";
break;
+ case PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR:
+ clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_HAS_DEVICE_ADDR,
+ clauses);
+ c_name = "has_device_addr";
+ break;
case PRAGMA_OMP_CLAUSE_IF:
clauses = cp_parser_omp_clause_if (parser, clauses, token->location,
true);
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_ALLOCATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IN_REDUCTION) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_THREAD_LIMIT) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR))
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR)\
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR))
static bool
cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
return false;
}
OMP_CLAUSE_DECL (c) = first;
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
+ return false;
OMP_CLAUSE_SIZE (c) = size;
if (TREE_CODE (t) == FIELD_DECL)
t = finish_non_static_data_member (t, NULL_TREE, NULL_TREE);
{
bitmap_head generic_head, firstprivate_head, lastprivate_head;
bitmap_head aligned_head, map_head, map_field_head, map_firstprivate_head;
- bitmap_head oacc_reduction_head;
+ bitmap_head oacc_reduction_head, is_on_device_head;
tree c, t, *pc;
tree safelen = NULL_TREE;
bool branch_seen = false;
/* If ort == C_ORT_OMP used as nontemporal_head or use_device_xxx_head
instead and for ort == C_ORT_OMP_TARGET used as in_reduction_head. */
bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack);
+ bitmap_initialize (&is_on_device_head, &bitmap_default_obstack);
if (ort & C_ORT_ACC)
for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
"%qD appears more than once in data clauses", t);
remove = true;
}
- else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_PRIVATE
+ else if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_PRIVATE
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR)
&& bitmap_bit_p (&map_head, DECL_UID (t)))
{
if (ort == C_ORT_ACC)
"%qD appears more than once in data clauses", t);
remove = true;
}
- else if (bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
+ 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)
error_at (OMP_CLAUSE_LOCATION (c),
case OMP_CLAUSE_USE_DEVICE_PTR:
field_ok = (ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP;
t = OMP_CLAUSE_DECL (c);
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR)
+ bitmap_set_bit (&is_on_device_head, DECL_UID (t));
if (!type_dependent_expression_p (t))
{
tree type = TREE_TYPE (t);
}
goto check_dup_generic;
+ case OMP_CLAUSE_HAS_DEVICE_ADDR:
+ t = OMP_CLAUSE_DECL (c);
+ if (TREE_CODE (t) == TREE_LIST)
+ {
+ if (handle_omp_array_sections (c, ort))
+ remove = true;
+ else
+ {
+ t = OMP_CLAUSE_DECL (c);
+ while (TREE_CODE (t) == INDIRECT_REF
+ || TREE_CODE (t) == ARRAY_REF)
+ t = TREE_OPERAND (t, 0);
+ }
+ }
+ bitmap_set_bit (&is_on_device_head, DECL_UID (t));
+ if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
+ cxx_mark_addressable (t);
+ goto check_dup_generic_t;
+
case OMP_CLAUSE_USE_DEVICE_ADDR:
field_ok = true;
t = OMP_CLAUSE_DECL (c);
case OMP_LIST_CACHE: type = "CACHE"; break;
case OMP_LIST_IS_DEVICE_PTR: type = "IS_DEVICE_PTR"; break;
case OMP_LIST_USE_DEVICE_PTR: type = "USE_DEVICE_PTR"; break;
+ case OMP_LIST_HAS_DEVICE_ADDR: type = "HAS_DEVICE_ADDR"; break;
case OMP_LIST_USE_DEVICE_ADDR: type = "USE_DEVICE_ADDR"; break;
case OMP_LIST_NONTEMPORAL: type = "NONTEMPORAL"; break;
case OMP_LIST_ALLOCATE: type = "ALLOCATE"; break;
OMP_LIST_USE_DEVICE_ADDR,
OMP_LIST_NONTEMPORAL,
OMP_LIST_ALLOCATE,
- OMP_LIST_NUM
+ OMP_LIST_HAS_DEVICE_ADDR,
+ OMP_LIST_NUM /* Must be the last. */
};
/* Because a symbol can belong to multiple namelists, they must be
OMP_MASK1_LAST
};
-/* OpenACC 2.0+ specific clauses. */
+/* More OpenMP clauses and OpenACC 2.0+ specific clauses. */
enum omp_mask2
{
OMP_CLAUSE_ASYNC,
OMP_CLAUSE_FINALIZE,
OMP_CLAUSE_ATTACH,
OMP_CLAUSE_NOHOST,
+ OMP_CLAUSE_HAS_DEVICE_ADDR, /* OpenMP 5.1 */
/* This must come last. */
OMP_MASK2_LAST
};
}
break;
case 'h':
+ if ((mask & OMP_CLAUSE_HAS_DEVICE_ADDR)
+ && gfc_match_omp_variable_list
+ ("has_device_addr (", &c->lists[OMP_LIST_HAS_DEVICE_ADDR],
+ false, NULL, NULL, true) == MATCH_YES)
+ continue;
if ((mask & OMP_CLAUSE_HINT)
&& (m = gfc_match_dupl_check (!c->hint, "hint", true, &c->hint))
!= MATCH_NO)
continue;
if ((mask & OMP_CLAUSE_USE_DEVICE_ADDR)
&& gfc_match_omp_variable_list
- ("use_device_addr (",
- &c->lists[OMP_LIST_USE_DEVICE_ADDR], false) == MATCH_YES)
+ ("use_device_addr (", &c->lists[OMP_LIST_USE_DEVICE_ADDR],
+ false, NULL, NULL, true) == MATCH_YES)
continue;
break;
case 'v':
| OMP_CLAUSE_DEPEND | OMP_CLAUSE_NOWAIT | OMP_CLAUSE_PRIVATE \
| OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULTMAP \
| OMP_CLAUSE_IS_DEVICE_PTR | OMP_CLAUSE_IN_REDUCTION \
- | OMP_CLAUSE_THREAD_LIMIT | OMP_CLAUSE_ALLOCATE)
+ | OMP_CLAUSE_THREAD_LIMIT | OMP_CLAUSE_ALLOCATE \
+ | OMP_CLAUSE_HAS_DEVICE_ADDR)
#define OMP_TARGET_DATA_CLAUSES \
(omp_mask (OMP_CLAUSE_DEVICE) | OMP_CLAUSE_MAP | OMP_CLAUSE_IF \
| OMP_CLAUSE_USE_DEVICE_PTR | OMP_CLAUSE_USE_DEVICE_ADDR)
"IN_REDUCTION", "TASK_REDUCTION",
"DEVICE_RESIDENT", "LINK", "USE_DEVICE",
"CACHE", "IS_DEVICE_PTR", "USE_DEVICE_PTR", "USE_DEVICE_ADDR",
- "NONTEMPORAL", "ALLOCATE" };
+ "NONTEMPORAL", "ALLOCATE", "HAS_DEVICE_ADDR" };
STATIC_ASSERT (ARRAY_SIZE (clause_names) == OMP_LIST_NUM);
if (omp_clauses == NULL)
n->sym->name, name, &n->where);
}
break;
+ case OMP_LIST_HAS_DEVICE_ADDR:
case OMP_LIST_USE_DEVICE_PTR:
case OMP_LIST_USE_DEVICE_ADDR:
/* FIXME: Handle OMP_LIST_USE_DEVICE_PTR. */
tree t = gfc_trans_omp_variable (namelist->sym, declare_simd);
if (t != error_mark_node)
{
- tree node = build_omp_clause (input_location, code);
+ tree node;
+ /* For HAS_DEVICE_ADDR of an array descriptor, firstprivatize the
+ descriptor such that the bounds are available; its data component
+ is unmodified; it is handled as device address inside target. */
+ if (code == OMP_CLAUSE_HAS_DEVICE_ADDR
+ && (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (t))
+ || (POINTER_TYPE_P (TREE_TYPE (t))
+ && GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (TREE_TYPE (t))))))
+ node = build_omp_clause (input_location, OMP_CLAUSE_FIRSTPRIVATE);
+ else
+ node = build_omp_clause (input_location, code);
OMP_CLAUSE_DECL (node) = t;
list = gfc_trans_add_clause (node, list);
case OMP_LIST_IS_DEVICE_PTR:
clause_code = OMP_CLAUSE_IS_DEVICE_PTR;
goto add_clause;
+ case OMP_LIST_HAS_DEVICE_ADDR:
+ clause_code = OMP_CLAUSE_HAS_DEVICE_ADDR;
+ goto add_clause;
case OMP_LIST_NONTEMPORAL:
clause_code = OMP_CLAUSE_NONTEMPORAL;
goto add_clause;
flags = GOVD_EXPLICIT;
goto do_add;
+ case OMP_CLAUSE_HAS_DEVICE_ADDR:
+ decl = OMP_CLAUSE_DECL (c);
+ while (TREE_CODE (decl) == INDIRECT_REF
+ || TREE_CODE (decl) == ARRAY_REF)
+ decl = TREE_OPERAND (decl, 0);
+ flags = GOVD_EXPLICIT;
+ goto do_add_decl;
+
case OMP_CLAUSE_IS_DEVICE_PTR:
flags = GOVD_FIRSTPRIVATE | GOVD_EXPLICIT;
goto do_add;
}
break;
+ case OMP_CLAUSE_HAS_DEVICE_ADDR:
+ decl = OMP_CLAUSE_DECL (c);
+ while (TREE_CODE (decl) == INDIRECT_REF
+ || TREE_CODE (decl) == ARRAY_REF)
+ decl = TREE_OPERAND (decl, 0);
+ n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
+ remove = n == NULL || !(n->value & GOVD_SEEN);
+ break;
+
+ case OMP_CLAUSE_IS_DEVICE_PTR:
case OMP_CLAUSE_NONTEMPORAL:
decl = OMP_CLAUSE_DECL (c);
n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
case OMP_CLAUSE_DETACH:
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_USE_DEVICE_ADDR:
- case OMP_CLAUSE_IS_DEVICE_PTR:
case OMP_CLAUSE_ASYNC:
case OMP_CLAUSE_WAIT:
case OMP_CLAUSE_INDEPENDENT:
decl = OMP_CLAUSE_DECL (c);
do_private:
if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE
- || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR)
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
&& is_gimple_omp_offloaded (ctx->stmt))
{
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
by_ref = !omp_privatize_by_reference (decl);
install_var_field (decl, by_ref, 3, ctx);
}
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
+ {
+ if (TREE_CODE (decl) == INDIRECT_REF)
+ decl = TREE_OPERAND (decl, 0);
+ install_var_field (decl, true, 3, ctx);
+ }
else if (TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
- install_var_field (decl, true, 3, ctx);
+ install_var_field (decl, true, 3, ctx);
else
install_var_field (decl, false, 3, ctx);
}
install_var_local (decl, ctx);
break;
+ case OMP_CLAUSE_HAS_DEVICE_ADDR:
+ decl = OMP_CLAUSE_DECL (c);
+ while (TREE_CODE (decl) == INDIRECT_REF
+ || TREE_CODE (decl) == ARRAY_REF)
+ decl = TREE_OPERAND (decl, 0);
+ goto do_private;
+
case OMP_CLAUSE_IS_DEVICE_PTR:
decl = OMP_CLAUSE_DECL (c);
goto do_private;
case OMP_CLAUSE_FIRSTPRIVATE:
case OMP_CLAUSE_PRIVATE:
case OMP_CLAUSE_LINEAR:
+ case OMP_CLAUSE_HAS_DEVICE_ADDR:
case OMP_CLAUSE_IS_DEVICE_PTR:
decl = OMP_CLAUSE_DECL (c);
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
+ {
+ while (TREE_CODE (decl) == INDIRECT_REF
+ || TREE_CODE (decl) == ARRAY_REF)
+ decl = TREE_OPERAND (decl, 0);
+ }
+
if (is_variable_sized (decl))
{
if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE
- || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR)
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
&& is_gimple_omp_offloaded (ctx->stmt))
{
tree decl2 = DECL_VALUE_EXPR (decl);
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_USE_DEVICE_ADDR:
+ case OMP_CLAUSE_HAS_DEVICE_ADDR:
case OMP_CLAUSE_IS_DEVICE_PTR:
var = OMP_CLAUSE_DECL (c);
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
+ {
+ while (TREE_CODE (var) == INDIRECT_REF
+ || TREE_CODE (var) == ARRAY_REF)
+ var = TREE_OPERAND (var, 0);
+ }
map_cnt++;
if (is_variable_sized (var))
{
SET_DECL_VALUE_EXPR (new_var, x);
DECL_HAS_VALUE_EXPR_P (new_var) = 1;
}
- else if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR
+ else if (((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
&& !omp_privatize_by_reference (var)
&& !omp_is_allocatable_or_ptr (var)
&& !lang_hooks.decls.omp_array_data (var, true))
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_USE_DEVICE_ADDR:
+ case OMP_CLAUSE_HAS_DEVICE_ADDR:
case OMP_CLAUSE_IS_DEVICE_PTR:
ovar = OMP_CLAUSE_DECL (c);
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
+ {
+ while (TREE_CODE (ovar) == INDIRECT_REF
+ || TREE_CODE (ovar) == ARRAY_REF)
+ ovar = TREE_OPERAND (ovar, 0);
+ }
var = lookup_decl_in_outer_ctx (ovar, ctx);
if (lang_hooks.decls.omp_array_data (ovar, true))
{
- tkind = (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR
+ tkind = ((OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR
+ && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_HAS_DEVICE_ADDR)
? GOMP_MAP_USE_DEVICE_PTR : GOMP_MAP_FIRSTPRIVATE_INT);
x = build_sender_ref ((splay_tree_key) &DECL_NAME (ovar), ctx);
}
- else if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR)
+ else if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR
+ && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_HAS_DEVICE_ADDR)
{
tkind = GOMP_MAP_USE_DEVICE_PTR;
x = build_sender_ref ((splay_tree_key) &DECL_UID (ovar), ctx);
type = TREE_TYPE (ovar);
if (lang_hooks.decls.omp_array_data (ovar, true))
var = lang_hooks.decls.omp_array_data (ovar, false);
- else if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR
+ else if (((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
&& !omp_privatize_by_reference (ovar)
&& !omp_is_allocatable_or_ptr (ovar))
|| TREE_CODE (type) == ARRAY_TYPE)
if (POINTER_TYPE_P (type)
&& TREE_CODE (type) != ARRAY_TYPE
&& ((OMP_CLAUSE_CODE (c) != OMP_CLAUSE_USE_DEVICE_ADDR
+ && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_HAS_DEVICE_ADDR
&& !omp_is_allocatable_or_ptr (ovar))
|| (omp_privatize_by_reference (ovar)
&& omp_is_allocatable_or_ptr (ovar))))
break;
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_USE_DEVICE_ADDR:
+ case OMP_CLAUSE_HAS_DEVICE_ADDR:
case OMP_CLAUSE_IS_DEVICE_PTR:
tree new_var;
gimple_seq assign_body;
var = OMP_CLAUSE_DECL (c);
is_array_data = lang_hooks.decls.omp_array_data (var, true) != NULL;
- if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR)
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR
+ && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_HAS_DEVICE_ADDR)
x = build_sender_ref (is_array_data
? (splay_tree_key) &DECL_NAME (var)
: (splay_tree_key) &DECL_UID (var), ctx);
else
- x = build_receiver_ref (var, false, ctx);
+ {
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
+ {
+ while (TREE_CODE (var) == INDIRECT_REF
+ || TREE_CODE (var) == ARRAY_REF)
+ var = TREE_OPERAND (var, 0);
+ }
+ x = build_receiver_ref (var, false, ctx);
+ }
if (is_array_data)
{
gimple_seq_add_stmt (&assign_body,
gimple_build_assign (new_var, x));
}
- else if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR
+ else if (((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
&& !omp_privatize_by_reference (var)
&& !omp_is_allocatable_or_ptr (var))
|| TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
type = TREE_TYPE (type);
if (POINTER_TYPE_P (type)
&& TREE_CODE (type) != ARRAY_TYPE
- && (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_USE_DEVICE_ADDR
+ && ((OMP_CLAUSE_CODE (c) != OMP_CLAUSE_USE_DEVICE_ADDR
+ && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_HAS_DEVICE_ADDR)
|| (omp_privatize_by_reference (var)
&& omp_is_allocatable_or_ptr (var))))
{
gimple_build_assign (new_var, x));
}
tree present;
- present = (do_optional_check
+ present = ((do_optional_check
+ && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_HAS_DEVICE_ADDR)
? omp_check_optional_argument (OMP_CLAUSE_DECL (c), true)
: NULL_TREE);
if (present)
}
void
-bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
+bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int hda, int s,
int nte, int tl, int nth, int g, int nta, int fi, int pp, int *q, int *dd, int ntm)
{
#pragma omp for simd \
#pragma omp target parallel \
device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \
if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) \
- nowait depend(inout: dd[0]) allocate (omp_default_mem_alloc:f) in_reduction(+:r2)
+ nowait depend(inout: dd[0]) allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr(hda)
;
#pragma omp target parallel for \
device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \
if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) \
lastprivate (l) linear (ll:1) ordered schedule(static, 4) collapse(1) nowait depend(inout: dd[0]) \
- allocate (omp_default_mem_alloc:f) in_reduction(+:r2)
+ allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr(hda)
for (int i = 0; i < 64; i++)
ll++;
#pragma omp target parallel for \
device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \
if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) \
lastprivate (l) linear (ll:1) schedule(static, 4) collapse(1) nowait depend(inout: dd[0]) order(concurrent) \
- allocate (omp_default_mem_alloc:f) in_reduction(+:r2)
+ allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr(hda)
for (int i = 0; i < 64; i++)
ll++;
#pragma omp target parallel for simd \
if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) \
lastprivate (l) linear (ll:1) schedule(static, 4) collapse(1) \
safelen(8) simdlen(4) aligned(q: 32) nowait depend(inout: dd[0]) nontemporal(ntm) if (simd: i3) order(concurrent) \
- allocate (omp_default_mem_alloc:f) in_reduction(+:r2)
+ allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr(hda)
for (int i = 0; i < 64; i++)
ll++;
#pragma omp target teams \
device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \
shared(s) default(shared) reduction(+:r) num_teams(nte - 1:nte) thread_limit(tl) nowait depend(inout: dd[0]) \
- allocate (omp_default_mem_alloc:f) in_reduction(+:r2)
+ allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr(hda)
;
#pragma omp target teams distribute \
device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \
shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) order(concurrent) \
- collapse(1) dist_schedule(static, 16) nowait depend(inout: dd[0]) allocate (omp_default_mem_alloc:f) in_reduction(+:r2)
+ collapse(1) dist_schedule(static, 16) nowait depend(inout: dd[0]) allocate (omp_default_mem_alloc:f) in_reduction(+:r2) \
+ has_device_addr(hda)
for (int i = 0; i < 64; i++)
;
#pragma omp target teams distribute parallel for \
collapse(1) dist_schedule(static, 16) \
if (parallel: i2) num_threads (nth) proc_bind(spread) \
lastprivate (l) schedule(static, 4) nowait depend(inout: dd[0]) order(concurrent) \
- allocate (omp_default_mem_alloc:f) in_reduction(+:r2)
+ allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr(hda)
for (int i = 0; i < 64; i++)
ll++;
#pragma omp target teams distribute parallel for simd \
if (parallel: i2) num_threads (nth) proc_bind(spread) \
lastprivate (l) schedule(static, 4) order(concurrent) \
safelen(8) simdlen(4) aligned(q: 32) nowait depend(inout: dd[0]) nontemporal(ntm) if (simd: i3) \
- allocate (omp_default_mem_alloc:f) in_reduction(+:r2)
+ allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr(hda)
for (int i = 0; i < 64; i++)
ll++;
#pragma omp target teams distribute simd \
shared(s) default(shared) reduction(+:r) num_teams(nte-1:nte) thread_limit(tl) \
collapse(1) dist_schedule(static, 16) order(concurrent) \
safelen(8) simdlen(4) aligned(q: 32) nowait depend(inout: dd[0]) nontemporal(ntm) \
- allocate (omp_default_mem_alloc:f) in_reduction(+:r2)
+ allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr(hda)
for (int i = 0; i < 64; i++)
ll++;
#pragma omp target simd \
device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \
safelen(8) simdlen(4) lastprivate (l) linear(ll: 1) aligned(q: 32) reduction(+:r) \
nowait depend(inout: dd[0]) nontemporal(ntm) if(simd:i3) order(concurrent) \
- allocate (omp_default_mem_alloc:f) in_reduction(+:r2)
+ allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr(hda)
for (int i = 0; i < 64; i++)
ll++;
#pragma omp taskgroup task_reduction(+:r2) allocate (r2)
device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \
if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) \
nowait depend(inout: dd[0]) lastprivate (l) bind(parallel) order(concurrent) collapse(1) \
- allocate (omp_default_mem_alloc: f) in_reduction(+:r2)
+ allocate (omp_default_mem_alloc: f) in_reduction(+:r2) has_device_addr(hda)
for (l = 0; l < 64; ++l)
;
#pragma omp target parallel loop \
device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \
if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) \
nowait depend(inout: dd[0]) lastprivate (l) order(concurrent) collapse(1) \
- allocate (omp_default_mem_alloc: f) in_reduction(+:r2)
+ allocate (omp_default_mem_alloc: f) in_reduction(+:r2) has_device_addr(hda)
for (l = 0; l < 64; ++l)
;
#pragma omp target teams loop \
device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \
shared(s) default(shared) reduction(+:r) num_teams(nte-1:nte) thread_limit(tl) nowait depend(inout: dd[0]) \
lastprivate (l) bind(teams) collapse(1) \
- allocate (omp_default_mem_alloc: f) in_reduction(+:r2)
+ allocate (omp_default_mem_alloc: f) in_reduction(+:r2) has_device_addr(hda)
for (l = 0; l < 64; ++l)
;
#pragma omp target teams loop \
device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \
shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) nowait depend(inout: dd[0]) \
lastprivate (l) order(concurrent) collapse(1) \
- allocate (omp_default_mem_alloc: f) in_reduction(+:r2)
+ allocate (omp_default_mem_alloc: f) in_reduction(+:r2) has_device_addr(hda)
for (l = 0; l < 64; ++l)
;
}
--- /dev/null
+/* { dg-do compile } */
+
+void
+foo ()
+{
+ int * x;
+ #pragma omp target is_device_ptr(x) has_device_addr(x) /*{ dg-error "'x' appears more than once in data clauses" } */
+ ;
+ #pragma omp target has_device_addr(x) is_device_ptr(x) /* { dg-error "'x' appears more than once in data clauses" } */
+ ;
+
+ int y = 42;
+ #pragma omp target has_device_addr(y) has_device_addr(y) /* { dg-error "'y' appears more than once in data clauses" } */
+ ;
+
+ #pragma omp target private(y) has_device_addr(y) /*{ dg-error "'y' appears more than once in data clauses" } */
+ ;
+ #pragma omp target has_device_addr(y) private(y) /*{ dg-error "'y' appears more than once in data clauses" } */
+ ;
+ #pragma omp target firstprivate(y) has_device_addr(y) /*{ dg-error "'y' appears more than once in data clauses" } */
+ ;
+
+ #pragma omp target has_device_addr(y) map(y) /* { dg-error "'y' appears both in data and map clauses" } */
+ ;
+ #pragma omp target map(y) has_device_addr(y) /* { dg-error "'y' appears both in data and map clauses" } */
+ ;
+
+ int z[3] = { 2, 5, 7 };
+ #pragma omp target data map(z[:3]) use_device_addr(z)
+ #pragma omp target has_device_addr(z[1:])
+ ;
+
+ #pragma omp target data map(z[:3]) use_device_addr(z)
+ #pragma omp target has_device_addr(z[1])
+ ;
+
+ #pragma omp target data map(z[:3]) use_device_addr(z)
+ #pragma omp target has_device_addr(z[1:2])
+ ;
+
+ #pragma omp target data map(z[:3]) use_device_addr(z)
+ #pragma omp target has_device_addr(z[:2])
+ ;
+
+ int w[3][4];
+ #pragma omp target data map(w) use_device_addr(w)
+ #pragma omp target has_device_addr(w[1][2])
+ ;
+
+ #pragma omp target data map(w) use_device_addr(w)
+ #pragma omp target has_device_addr(w[:1][2:])
+ ;
+
+ int u[0];
+ #pragma omp target data map(u) use_device_addr(u)
+ #pragma omp target has_device_addr(u)
+ ;
+
+ struct S { int m; } s;
+ s.m = 42;
+ #pragma omp target data map (s) use_device_addr (s)
+ #pragma omp target has_device_addr (s)
+ ++s.m;
+
+}
--- /dev/null
+/* { dg-do compile } */
+/* { dg-options "-fopenmp -fdump-tree-gimple" } */
+
+void
+foo ()
+{
+ int x, y;
+
+ #pragma omp target data map(x, y) use_device_addr(x, y)
+ #pragma omp target has_device_addr(x, y)
+ {
+ x = 42;
+ }
+}
+
+/* { dg-final { scan-tree-dump "has_device_addr\\(x\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump-not "has_device_addr\\(y\\)" "gimple" } } */
--- /dev/null
+/* { dg-do compile } */
+
+void
+foo ()
+{
+ int *x;
+
+ #pragma omp target is_device_ptr(x) is_device_ptr(x) /* { dg-error "'x' appears more than once in data clauses" } */
+ ;
+
+ #pragma omp target private(x) is_device_ptr(x) /*{ dg-error "'x' appears more than once in data clauses" } */
+ ;
+ #pragma omp target is_device_ptr(x) private(x) /*{ dg-error "'x' appears more than once in data clauses" } */
+ ;
+ #pragma omp target firstprivate(x) is_device_ptr(x) /*{ dg-error "'x' appears more than once in data clauses" } */
+ ;
+
+ #pragma omp target is_device_ptr(x) map(x) /* { dg-error "'x' appears both in data and map clauses" } */
+ ;
+ #pragma omp target map(x) is_device_ptr(x) /* { dg-error "'x' appears both in data and map clauses" } */
+ ;
+}
--- /dev/null
+/* { dg-do compile } */
+/* { dg-options "-fopenmp -fdump-tree-gimple" } */
+
+void
+foo ()
+{
+ int *x, *y;
+
+ #pragma omp target data map(x, y) use_device_ptr(x, y)
+ #pragma omp target is_device_ptr(x, y)
+ {
+ *x = 42;
+ }
+}
+
+/* { dg-final { scan-tree-dump "is_device_ptr\\(x\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump-not "is_device_ptr\\(y\\)" "gimple" } } */
}
void
-bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
+bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int hda, int s,
int nte, int tl, int nth, int g, int nta, int fi, int pp, int *q, int *dd, int ntm,
const char *msg)
{
[[omp::directive (target parallel
device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp)
if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread)
- nowait depend(inout: dd[0]) allocate (omp_default_mem_alloc:f) in_reduction(+:r2))]]
+ nowait depend(inout: dd[0]) allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr (hda))]]
;
[[omp::directive (target parallel for
device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp)
if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread)
lastprivate (l) linear (ll:1) ordered schedule(static, 4) collapse(1) nowait depend(inout: dd[0])
- allocate (omp_default_mem_alloc:f) in_reduction(+:r2))]]
+ allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr (hda))]]
for (int i = 0; i < 64; i++)
ll++;
[[omp::directive (target parallel for
device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp)
if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread)
lastprivate (l) linear (ll:1) schedule(static, 4) collapse(1) nowait depend(inout: dd[0]) order(concurrent)
- allocate (omp_default_mem_alloc:f) in_reduction(+:r2))]]
+ allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr (hda))]]
for (int i = 0; i < 64; i++)
ll++;
[[omp::sequence (omp::directive (target parallel for simd
if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread)
lastprivate (l) linear (ll:1) schedule(static, 4) collapse(1)
safelen(8) simdlen(4) aligned(q: 32) nowait depend(inout: dd[0]) nontemporal(ntm) if (simd: i3) order(concurrent)
- allocate (omp_default_mem_alloc:f) in_reduction(+:r2)))]]
+ allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr (hda)))]]
for (int i = 0; i < 64; i++)
ll++;
[[omp::sequence (directive (target teams
device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp)
shared(s) default(shared) reduction(+:r) num_teams(nte-1:nte) thread_limit(tl) nowait depend(inout: dd[0])
- allocate (omp_default_mem_alloc:f) in_reduction(+:r2)))]]
+ allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr (hda)))]]
;
[[omp::sequence (directive (target
device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp)
- nowait depend(inout: dd[0]) allocate (omp_default_mem_alloc:f) in_reduction(+:r2)))]]
+ nowait depend(inout: dd[0]) allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr (hda)))]]
;
[[omp::sequence (omp::directive (target teams distribute
device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp)
shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) order(concurrent)
- collapse(1) dist_schedule(static, 16) nowait depend(inout: dd[0]) allocate (omp_default_mem_alloc:f) in_reduction(+:r2)))]]
+ collapse(1) dist_schedule(static, 16) nowait depend(inout: dd[0]) allocate (omp_default_mem_alloc:f) in_reduction(+:r2)
+ has_device_addr (hda)))]]
for (int i = 0; i < 64; i++)
;
[[omp::directive (target teams distribute parallel for
collapse(1) dist_schedule(static, 16)
if (parallel: i2) num_threads (nth) proc_bind(spread)
lastprivate (l) schedule(static, 4) nowait depend(inout: dd[0]) order(concurrent)
- allocate (omp_default_mem_alloc:f) in_reduction(+:r2))]]
+ allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr (hda))]]
for (int i = 0; i < 64; i++)
ll++;
[[omp::directive (target teams distribute parallel for simd
if (parallel: i2) num_threads (nth) proc_bind(spread)
lastprivate (l) schedule(static, 4) order(concurrent)
safelen(8) simdlen(4) aligned(q: 32) nowait depend(inout: dd[0]) nontemporal(ntm) if (simd: i3)
- allocate (omp_default_mem_alloc:f) in_reduction(+:r2))]]
+ allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr (hda))]]
for (int i = 0; i < 64; i++)
ll++;
[[omp::directive (target teams distribute simd
shared(s) default(shared) reduction(+:r) num_teams(nte-1:nte) thread_limit(tl)
collapse(1) dist_schedule(static, 16) order(concurrent)
safelen(8) simdlen(4) aligned(q: 32) nowait depend(inout: dd[0]) nontemporal(ntm)
- allocate (omp_default_mem_alloc:f) in_reduction(+:r2))]]
+ allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr (hda))]]
for (int i = 0; i < 64; i++)
ll++;
[[omp::directive (target simd
device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp)
safelen(8) simdlen(4) lastprivate (l) linear(ll: 1) aligned(q: 32) reduction(+:r)
nowait depend(inout: dd[0]) nontemporal(ntm) if(simd:i3) order(concurrent)
- allocate (omp_default_mem_alloc:f) in_reduction(+:r2))]]
+ allocate (omp_default_mem_alloc:f) in_reduction(+:r2) has_device_addr (hda))]]
for (int i = 0; i < 64; i++)
ll++;
[[omp::sequence (directive (taskgroup task_reduction(+:r2) allocate (r2)),
device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp)
if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread)
nowait depend(inout: dd[0]) lastprivate (l) bind(parallel) order(concurrent) collapse(1)
- allocate (omp_default_mem_alloc: f) in_reduction(+:r2))]]
+ allocate (omp_default_mem_alloc: f) in_reduction(+:r2) has_device_addr (hda))]]
for (l = 0; l < 64; ++l)
;
[[omp::directive (target parallel loop
device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp)
if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread)
nowait depend(inout: dd[0]) lastprivate (l) order(concurrent) collapse(1)
- allocate (omp_default_mem_alloc: f) in_reduction(+:r2))]]
+ allocate (omp_default_mem_alloc: f) in_reduction(+:r2) has_device_addr (hda))]]
for (l = 0; l < 64; ++l)
;
[[omp::directive (target teams loop
device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp)
shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) nowait depend(inout: dd[0])
lastprivate (l) bind(teams) collapse(1)
- allocate (omp_default_mem_alloc: f) in_reduction(+:r2))]]
+ allocate (omp_default_mem_alloc: f) in_reduction(+:r2) has_device_addr (hda))]]
for (l = 0; l < 64; ++l)
;
[[omp::directive (target teams loop
device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp)
shared(s) default(shared) reduction(+:r) num_teams(nte - 1 : nte) thread_limit(tl) nowait depend(inout: dd[0])
lastprivate (l) order(concurrent) collapse(1)
- allocate (omp_default_mem_alloc: f) in_reduction(+:r2))]]
+ allocate (omp_default_mem_alloc: f) in_reduction(+:r2) has_device_addr (hda))]]
for (l = 0; l < 64; ++l)
;
[[omp::directive (critical)]] {
}
void
-bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int s,
+bar (int d, int m, int i1, int i2, int i3, int p, int *idp, int hda, int s,
int nte, int tl, int nth, int g, int nta, int fi, int pp, int *q, int *dd, int ntm,
const char *msg)
{
[[omp::directive (target parallel,
device(d),map (tofrom: m),if (target: i1),private (p),firstprivate (f),defaultmap(tofrom: scalar),is_device_ptr (idp),
if (parallel: i2),default(shared),shared(s),reduction(+:r),num_threads (nth),proc_bind(spread)
- nowait depend(inout: dd[0]),allocate (omp_default_mem_alloc:f),in_reduction(+:r2))]]
+ nowait depend(inout: dd[0]),allocate (omp_default_mem_alloc:f),in_reduction(+:r2),has_device_addr (hda))]]
;
[[omp::directive (target parallel for,
device(d),map (tofrom: m),if (target: i1),private (p),firstprivate (f),defaultmap(tofrom: scalar),is_device_ptr (idp),
if (parallel: i2),default(shared),shared(s),reduction(+:r),num_threads (nth),proc_bind(spread),
lastprivate (l),linear (ll:1),ordered schedule(static, 4),collapse(1),nowait depend(inout: dd[0]),
- allocate (omp_default_mem_alloc:f),in_reduction(+:r2))]]
+ allocate (omp_default_mem_alloc:f),in_reduction(+:r2),has_device_addr (hda))]]
for (int i = 0; i < 64; i++)
ll++;
[[using omp:directive (target parallel for,
device(d),map (tofrom: m),if (target: i1),private (p),firstprivate (f),defaultmap(tofrom: scalar),is_device_ptr (idp),
if (parallel: i2),default(shared),shared(s),reduction(+:r),num_threads (nth),proc_bind(spread),
lastprivate (l),linear (ll:1),schedule(static, 4),collapse(1),nowait depend(inout: dd[0]),order(concurrent),
- allocate (omp_default_mem_alloc:f),in_reduction(+:r2))]]
+ allocate (omp_default_mem_alloc:f),in_reduction(+:r2),has_device_addr (hda))]]
for (int i = 0; i < 64; i++)
ll++;
[[omp::sequence (omp::directive (target parallel for simd,
if (parallel: i2),default(shared),shared(s),reduction(+:r),num_threads (nth),proc_bind(spread),
lastprivate (l),linear (ll:1),schedule(static, 4),collapse(1),
safelen(8),simdlen(4),aligned(q: 32),nowait depend(inout: dd[0]),nontemporal(ntm),if (simd: i3),order(concurrent),
- allocate (omp_default_mem_alloc:f),in_reduction(+:r2)))]]
+ allocate (omp_default_mem_alloc:f),in_reduction(+:r2),has_device_addr (hda)))]]
for (int i = 0; i < 64; i++)
ll++;
[[using omp:sequence (directive (target teams,
device(d),map (tofrom: m),if (target: i1),private (p),firstprivate (f),defaultmap(tofrom: scalar),is_device_ptr (idp),
- shared(s),default(shared),reduction(+:r),num_teams(nte),thread_limit(tl),nowait, depend(inout: dd[0]),
- allocate (omp_default_mem_alloc:f) in_reduction(+:r2)))]]
+ shared(s),default(shared),reduction(+:r),num_teams(nte),thread_limit(tl),nowait,depend(inout: dd[0]),
+ allocate (omp_default_mem_alloc:f),in_reduction(+:r2),has_device_addr (hda)))]]
;
[[using omp:sequence (directive (target,
device(d),map (tofrom: m),if (target: i1),private (p),firstprivate (f),defaultmap(tofrom: scalar),is_device_ptr (idp),
- nowait depend(inout: dd[0]),allocate (omp_default_mem_alloc:f),in_reduction(+:r2)))]]
+ nowait depend(inout: dd[0]),allocate (omp_default_mem_alloc:f),in_reduction(+:r2),has_device_addr(hda)))]]
;
[[omp::sequence (omp::directive (target teams distribute,
device(d),map (tofrom: m),if (target: i1),private (p),firstprivate (f),defaultmap(tofrom: scalar),is_device_ptr (idp),
shared(s),default(shared),reduction(+:r),num_teams(nte-1:nte),thread_limit(tl),order(concurrent),
- collapse(1),dist_schedule(static, 16),nowait depend(inout: dd[0]),allocate (omp_default_mem_alloc:f),in_reduction(+:r2)))]]
+ collapse(1),dist_schedule(static, 16),nowait depend(inout: dd[0]),allocate (omp_default_mem_alloc:f),in_reduction(+:r2),
+ has_device_addr (hda)))]]
for (int i = 0; i < 64; i++)
;
[[omp::directive (target teams distribute parallel for,
collapse(1),dist_schedule(static, 16),
if (parallel: i2),num_threads (nth),proc_bind(spread),
lastprivate (l),schedule(static, 4),nowait depend(inout: dd[0]),order(concurrent),
- allocate (omp_default_mem_alloc:f),in_reduction(+:r2))]]
+ allocate (omp_default_mem_alloc:f),in_reduction(+:r2),has_device_addr (hda))]]
for (int i = 0; i < 64; i++)
ll++;
[[omp::directive (target teams distribute parallel for simd,
if (parallel: i2),num_threads (nth),proc_bind(spread),
lastprivate (l),schedule(static, 4),order(concurrent),
safelen(8),simdlen(4),aligned(q: 32),nowait depend(inout: dd[0]),nontemporal(ntm),if (simd: i3),
- allocate (omp_default_mem_alloc:f),in_reduction(+:r2))]]
+ allocate (omp_default_mem_alloc:f),in_reduction(+:r2),has_device_addr (hda))]]
for (int i = 0; i < 64; i++)
ll++;
[[omp::directive (target teams distribute simd,
shared(s),default(shared),reduction(+:r),num_teams(nte),thread_limit(tl),
collapse(1),dist_schedule(static, 16),order(concurrent),
safelen(8),simdlen(4),aligned(q: 32),nowait depend(inout: dd[0]),nontemporal(ntm),
- allocate (omp_default_mem_alloc:f),in_reduction(+:r2))]]
+ allocate (omp_default_mem_alloc:f),in_reduction(+:r2),has_device_addr (hda))]]
for (int i = 0; i < 64; i++)
ll++;
[[omp::directive (target simd,
device(d),map (tofrom: m),if (target: i1),private (p),firstprivate (f),defaultmap(tofrom: scalar),is_device_ptr (idp),
safelen(8),simdlen(4),lastprivate (l),linear(ll: 1),aligned(q: 32),reduction(+:r),
nowait depend(inout: dd[0]),nontemporal(ntm),if(simd:i3),order(concurrent),
- allocate (omp_default_mem_alloc:f),in_reduction(+:r2))]]
+ allocate (omp_default_mem_alloc:f),in_reduction(+:r2),has_device_addr (hda))]]
for (int i = 0; i < 64; i++)
ll++;
[[omp::sequence (directive (taskgroup, task_reduction(+:r2), allocate (r2)),
device(d),map (tofrom: m),if (target: i1),private (p),firstprivate (f),defaultmap(tofrom: scalar),is_device_ptr (idp),
if (parallel: i2),default(shared),shared(s),reduction(+:r),num_threads (nth),proc_bind(spread),
nowait depend(inout: dd[0]),lastprivate (l),bind(parallel),order(concurrent),collapse(1),
- allocate (omp_default_mem_alloc: f),in_reduction(+:r2))]]
+ allocate (omp_default_mem_alloc: f),in_reduction(+:r2),has_device_addr (hda))]]
for (l = 0; l < 64; ++l)
;
[[omp::directive (target parallel loop,
device(d),map (tofrom: m),if (target: i1),private (p),firstprivate (f),defaultmap(tofrom: scalar),is_device_ptr (idp),
if (parallel: i2),default(shared),shared(s),reduction(+:r),num_threads (nth),proc_bind(spread),
nowait depend(inout: dd[0]),lastprivate (l),order(concurrent),collapse(1),
- allocate (omp_default_mem_alloc: f),in_reduction(+:r2))]]
+ allocate (omp_default_mem_alloc: f),in_reduction(+:r2),has_device_addr (hda))]]
for (l = 0; l < 64; ++l)
;
[[omp::directive (target teams loop,
device(d),map (tofrom: m),if (target: i1),private (p),firstprivate (f),defaultmap(tofrom: scalar),is_device_ptr (idp),
shared(s),default(shared),reduction(+:r),num_teams(nte-1:nte),thread_limit(tl),nowait,depend(inout: dd[0]),
lastprivate (l),bind(teams),collapse(1),
- allocate (omp_default_mem_alloc: f),in_reduction(+:r2))]]
+ allocate (omp_default_mem_alloc: f),in_reduction(+:r2),has_device_addr (hda))]]
for (l = 0; l < 64; ++l)
;
[[omp::directive (target teams loop,
device(d),map (tofrom: m),if (target: i1),private (p),firstprivate (f),defaultmap(tofrom: scalar),is_device_ptr (idp),
shared(s),default(shared),reduction(+:r),num_teams(nte),thread_limit(tl),nowait,depend(inout: dd[0]),
lastprivate (l),order(concurrent),collapse(1)
- allocate (omp_default_mem_alloc: f),in_reduction(+:r2))]]
+ allocate (omp_default_mem_alloc: f),in_reduction(+:r2),has_device_addr (hda))]]
for (l = 0; l < 64; ++l)
;
[[omp::directive (critical)]] {
--- /dev/null
+! Test to ensure that IS_DEVICE_PTR is removed for non-used variables.
+
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+program main
+ use iso_c_binding
+ implicit none
+
+ integer :: x, y
+ call foo (x, y)
+
+contains
+ subroutine foo (a, b)
+ integer, target :: a, b
+
+ !$omp target data map(a, b) use_device_ptr(a, b)
+ !$omp target is_device_ptr(a, b)
+ a = 42
+ !$omp end target
+ !$omp end target data
+ end subroutine foo
+
+end program main
+
+! { dg-final { scan-tree-dump "is_device_ptr\\(a\\)" "gimple" } }
+! { dg-final { scan-tree-dump-not "is_device_ptr\\(b\\)" "gimple" } }
--- /dev/null
+! { dg-do compile }
+
+implicit none
+
+integer, target :: x
+integer, pointer :: ptr
+integer :: a(5)
+
+!$omp target has_device_addr(x)
+!$omp end target
+!$omp target has_device_addr(ptr)
+!$omp end target
+!$omp target has_device_addr(a)
+!$omp end target
+!$omp target has_device_addr(a(2:3))
+!$omp end target
+!$omp target has_device_addr(a(:3))
+!$omp end target
+!$omp target has_device_addr(a(2:))
+!$omp end target
+!$omp target has_device_addr(a(2))
+!$omp end target
+
+!$omp target has_device_addr(x) has_device_addr(x) ! { dg-error "'x' present on multiple clauses" }
+!$omp end target
+
+!$omp target private(x) has_device_addr(x) ! { dg-error "'x' present on multiple clauses" }
+!$omp end target
+!$omp target has_device_addr(x) private(x) ! { dg-error "'x' present on multiple clauses" }
+!$omp end target
+!$omp target firstprivate(x) has_device_addr(x) ! { dg-error "'x' present on multiple clauses" }
+!$omp end target
+!$omp target has_device_addr(x) firstprivate(x) ! { dg-error "'x' present on multiple clauses" }
+!$omp end target
+
+end
\ No newline at end of file
--- /dev/null
+! Test to ensure that HAS_DEVICE_ADDR is removed for non-used variables.
+
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+program main
+ use iso_c_binding
+ implicit none
+
+ integer :: x, y
+ call foo (x, y)
+
+contains
+ subroutine foo (a, b)
+ integer :: a, b
+
+ !$omp target data map(a) use_device_addr(a)
+ !$omp target has_device_addr(a)
+ a = 42
+ !$omp end target
+ !$omp end target data
+ end subroutine foo
+
+end program main
+
+! { dg-final { scan-tree-dump "has_device_addr\\(a\\)" "gimple" } }
+! { dg-final { scan-tree-dump-not "has_device_addr\\(b\\)" "gimple" } }
OpenMP clause: map ({alloc:,to:,from:,tofrom:,}variable-list). */
OMP_CLAUSE_MAP,
+ /* OpenMP clause: has_device_addr (variable-list). */
+ OMP_CLAUSE_HAS_DEVICE_ADDR,
+
/* Internal structure to hold OpenACC cache directive's variable-list.
#pragma acc cache (variable-list). */
OMP_CLAUSE__CACHE_,
case OMP_CLAUSE_LINK:
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_USE_DEVICE_ADDR:
+ case OMP_CLAUSE_HAS_DEVICE_ADDR:
case OMP_CLAUSE_IS_DEVICE_PTR:
case OMP_CLAUSE_DETACH:
do_decl_clause:
case OMP_CLAUSE_LINK:
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_USE_DEVICE_ADDR:
+ case OMP_CLAUSE_HAS_DEVICE_ADDR:
case OMP_CLAUSE_IS_DEVICE_PTR:
case OMP_CLAUSE_DETACH:
do_decl_clause:
case OMP_CLAUSE_USE_DEVICE_ADDR:
name = "use_device_addr";
goto print_remap;
+ case OMP_CLAUSE_HAS_DEVICE_ADDR:
+ name = "has_device_addr";
+ goto print_remap;
case OMP_CLAUSE_IS_DEVICE_PTR:
name = "is_device_ptr";
goto print_remap;
2, /* OMP_CLAUSE_FROM */
2, /* OMP_CLAUSE_TO */
2, /* OMP_CLAUSE_MAP */
+ 1, /* OMP_CLAUSE_HAS_DEVICE_ADDR */
2, /* OMP_CLAUSE__CACHE_ */
2, /* OMP_CLAUSE_GANG */
1, /* OMP_CLAUSE_ASYNC */
"from",
"to",
"map",
+ "has_device_addr",
"_cache_",
"gang",
"async",
@item @code{align} clause/modifier in @code{allocate} directive/clause
and @code{allocator} directive @tab P @tab C/C++ on clause only
@item @code{thread_limit} clause to @code{target} construct @tab Y @tab
-@item @code{has_device_addr} clause to @code{target} construct @tab N @tab
+@item @code{has_device_addr} clause to @code{target} construct @tab Y @tab
@item iterators in @code{target update} motion clauses and @code{map}
clauses @tab N @tab
@item indirect calls to the device version of a procedure or function in
tgt_size = 0;
size_t i;
for (i = 0; i < mapnum; i++)
- if ((kinds[i] & 0xff) == GOMP_MAP_FIRSTPRIVATE)
+ if ((kinds[i] & 0xff) == GOMP_MAP_FIRSTPRIVATE && hostaddrs[i] != NULL)
{
size_t align = (size_t) 1 << (kinds[i] >> 8);
tgt_size = (tgt_size + align - 1) & ~(align - 1);
--- /dev/null
+/* Testing 'has_device_addr' clause on the target construct with reference. */
+
+#include <omp.h>
+
+int
+main ()
+{
+ int *dp = (int*)omp_target_alloc (sizeof(int), 0);
+
+ #pragma omp target is_device_ptr(dp)
+ *dp = 42;
+
+ int &x = *dp;
+
+ #pragma omp target has_device_addr(x)
+ x = 24;
+
+ #pragma omp target has_device_addr(x)
+ if (x != 24)
+ __builtin_abort ();
+
+ omp_target_free(dp, 0);
+}
--- /dev/null
+#include <omp.h>
+
+int
+main ()
+{
+ int *dp = (int*)omp_target_alloc (30*sizeof(int), 0);
+
+ #pragma omp target is_device_ptr(dp)
+ for (int i = 0; i < 30; i++)
+ dp[i] = i;
+
+ int (&x)[30] = *static_cast<int(*)[30]>(static_cast<void*>(dp));
+
+ #pragma omp target has_device_addr(x)
+ for (int i = 0; i < 30; i++)
+ x[i] = 2 * i;
+
+ #pragma omp target has_device_addr(x)
+ for (int i = 0; i < 30; i++)
+ if (x[i] != 2 * i)
+ __builtin_abort ();
+
+ #pragma omp target has_device_addr(x[1:5])
+ for (int i = 1; i < 6; i++)
+ x[i] = 3 * i;
+
+ #pragma omp target has_device_addr(x[1:5])
+ for (int i = 1; i < 6; i++)
+ if (x[i] != 3 * i)
+ __builtin_abort ();
+
+ omp_target_free (dp, 0);
+}
--- /dev/null
+/* Testing 'has_device_addr' clause on the target construct with reference. */
+
+#include <omp.h>
+
+int
+main ()
+{
+ int *dpx = (int*)omp_target_alloc (sizeof(int), 0);
+ int **dpy = (int**)omp_target_alloc (sizeof(int*), 0);
+
+ #pragma omp target is_device_ptr(dpx, dpy)
+ {
+ *dpx = 42;
+ int z = 77;
+ *dpy = &z;
+ }
+
+ int& x = *dpx;
+ int*& y = *dpy;
+
+ #pragma omp target has_device_addr(x, y)
+ {
+ x = 24;
+ y = &x;
+ }
+
+ #pragma omp target has_device_addr(x, y)
+ if (x != 24 || y != &x)
+ __builtin_abort ();
+
+ omp_target_free(dpx, 0);
+ omp_target_free(dpy, 0);
+}
--- /dev/null
+/* Testing 'has_device_addr' clause on the target construct with reference. */
+
+#include <omp.h>
+
+int
+main ()
+{
+ int *dpx = (int*)omp_target_alloc (sizeof(int), 0);
+ double *dpy = (double*)omp_target_alloc (sizeof(double), 0);
+
+ #pragma omp target is_device_ptr(dpx, dpy)
+ {
+ *dpx = 42;
+ *dpy = 43.5;
+ }
+
+ int &x = *dpx;
+ double &y = *dpy;
+
+ #pragma omp target has_device_addr(x, y)
+ {
+ x = 24;
+ y = 25.7;
+ }
+
+ #pragma omp target has_device_addr(y, x)
+ if (x != 24 || y != 25.7)
+ __builtin_abort ();
+
+ omp_target_free(dpx, 0);
+ omp_target_free(dpy, 0);
+}
--- /dev/null
+/* Testing the 'has_device_addr' clause on the target construct with
+ enclosing 'target data' construct. */
+
+#define N 40
+
+int
+main ()
+{
+ int x = 24;
+
+ #pragma omp target data map(x) use_device_addr(x)
+ #pragma omp target has_device_addr(x)
+ x = 42;
+ if (x != 42)
+ __builtin_abort ();
+
+ int y[N];
+
+ for (int i = 0; i < N; i++)
+ y[i] = 42;
+ #pragma omp target data map(y) use_device_addr(y)
+ #pragma omp target has_device_addr(y)
+ for (int i = 0; i < N; i++)
+ y[i] = i;
+ for (int i = 0; i < N; i++)
+ if (y[i] != i)
+ __builtin_abort ();
+
+ #pragma omp target data map(y[:N]) use_device_addr(y)
+ #pragma omp target has_device_addr(y[:N])
+ for (int i = 0; i < N; i++)
+ y[i] = i + 2;
+ for (int i = 0; i < N; i++)
+ if (y[i] != i + 2)
+ __builtin_abort ();
+
+ #pragma omp target data map(y[:N]) use_device_addr(y)
+ #pragma omp target has_device_addr(y[24])
+ y[24] = 42;
+ if (y[24] != 42)
+ __builtin_abort ();
+
+ #pragma omp target data map(y[:N]) use_device_addr(y)
+ #pragma omp target has_device_addr(y[24:])
+ for (int i = 24; i < N; i++)
+ y[i] = i + 3;
+ for (int i = 24; i < N; i++)
+ if (y[i] != i + 3)
+ __builtin_abort ();
+
+ #pragma omp target data map(y[:N]) use_device_addr(y)
+ #pragma omp target has_device_addr(y[12:24])
+ for (int i = 12; i < 24; i++)
+ y[i] = i + 4;
+ for (int i = 12; i < 24; i++)
+ if (y[i] != i + 4)
+ __builtin_abort ();
+
+ int u[0];
+ #pragma omp target data map(u) use_device_addr(u)
+ #pragma omp target has_device_addr(u)
+ ;
+
+ struct S { int m; } s;
+ s.m = 42;
+ #pragma omp target data map (s) use_device_addr (s)
+ #pragma omp target has_device_addr (s)
+ ++s.m;
+ if (s.m != 43)
+ __builtin_abort ();
+
+ return 0;
+}
--- /dev/null
+/* Testing 'has_device_addr' clause with variable sized array. */
+
+int
+foo (int size)
+{
+ int x[size];
+
+ #pragma omp target data map(x[:size]) use_device_addr(x)
+ #pragma omp target has_device_addr(x)
+ for (int i = 0; i < size; i++)
+ x[i] = i;
+ for (int i = 0; i < size; i++)
+ if (x[i] != i)
+ __builtin_abort ();
+
+ #pragma omp target data map(x) use_device_addr(x)
+ #pragma omp target has_device_addr(x[2:3])
+ for (int i = 0; i < size; i++)
+ x[i] = i;
+ for (int i = 0; i < size; i++)
+ if (x[i] != i)
+ __builtin_abort ();
+
+ return 0;
+}
+
+int
+main ()
+{
+ foo (40);
+
+ return 0;
+}
--- /dev/null
+program main
+ use omp_lib
+ use iso_c_binding
+ implicit none
+
+ integer, parameter :: N = 40
+ integer :: x, i
+ integer :: y (N)
+ integer :: u (0)
+
+ x = 24
+ !$omp target data map(x) use_device_addr(x)
+ !$omp target has_device_addr(x)
+ x = 42;
+ !$omp end target
+ !$omp end target data
+ if (x /= 42) stop 1
+
+ y = 42
+ !$omp target data map(y) use_device_addr(y)
+ !$omp target has_device_addr(y)
+ y = [(i, i=1, N)]
+ !$omp end target
+ !$omp end target data
+ if (any (y /= [(i, i = 1, N)])) stop 2
+
+ !$omp target data map(y(:N)) use_device_addr(y)
+ !$omp target has_device_addr(y(:N))
+ y = [(i+2, i=1, N)]
+ !$omp end target
+ !$omp end target data
+ if (any (y /= [(i+2, i = 1, N)])) stop 3
+
+ !$omp target data map(y) use_device_addr(y)
+ !$omp target has_device_addr(y(24:))
+ do i = 24, N
+ y(i) = i + 3
+ end do
+ !$omp end target
+ !$omp end target data
+ do i = 24, N
+ if (y(i) /= i + 3) stop 5
+ end do
+
+ !$omp target data map(u) use_device_addr(u)
+ !$omp target has_device_addr(u)
+ !$omp end target
+ !$omp end target data
+
+end program main
--- /dev/null
+program main
+ use omp_lib
+ use iso_c_binding
+ implicit none
+
+ integer, parameter :: N = 5
+ integer :: i, x(N), y(N), z(N:2*N-1)
+ target :: z
+
+ x = 42
+ y = 43
+ z = 44
+
+ call foo (x, y, z)
+ if (any (x /= [(i, i = 1, N)])) stop 1
+ if (any (y /= [(2*i, i = 1, N)])) stop 2
+ if (any (z /= [(3*i, i = 1, N)])) stop 3
+
+ contains
+ subroutine foo(a, b, c)
+ integer :: a(:)
+ integer :: b(*)
+ integer, pointer, intent(in) :: c(:)
+
+ !$omp target data map(a,b(:N),c) use_device_addr(a,b(:N),c)
+ !$omp target has_device_addr(A,B(:N),C)
+ if (lbound(a,dim=1) /= 1 .or. ubound(a,dim=1) /= N) stop 10
+ if (lbound(b,dim=1) /= 1) stop 11
+ if (lbound(c,dim=1) /= N .or. ubound(c,dim=1) /= 2*N-1) stop 12
+ if (any (a /= 42)) stop 13
+ if (any (b(:N) /= 43)) stop 14
+ if (any (c /= 44)) stop 15
+ a = [(i, i=1, N)]
+ b(:N) = [(2*i, i = 1, N)]
+ c = [(3*i, i = 1, N)]
+ !$omp end target
+ !$omp end target data
+ end subroutine foo
+
+end program main
--- /dev/null
+! Test optional dummy arguments in HAS_DEVICE_ADDR.
+
+program main
+ use omp_lib
+ use iso_c_binding
+ implicit none
+
+ integer, target :: x
+ integer, pointer :: ptr
+ integer, parameter :: N=7
+ real :: y1(N), y2(N)
+ integer, target :: y3(N:2*N-1)
+ integer :: i
+
+ x = 24
+ ptr => x
+ y1 = 42.24
+ y2 = 42.24
+ y3 = 42
+
+ call optional_scalar (is_present=.false.)
+ if (x /= 24) stop 1
+
+ call optional_scalar (x, is_present=.true.)
+ if (x /= 42) stop 2
+
+ call optional_ptr (is_present=.false.)
+ if (x /= 42) stop 3
+ if (ptr /= 42) stop 4
+
+ call optional_ptr (ptr, is_present=.true.)
+ if (x /= 84) stop 5
+ if (ptr /= 84) stop 6
+
+ call optional_array (is_present=.false.)
+ if (any (y1 /= [(42.24, i=1, N)])) stop 7
+ if (any (y2 /= [(42.24, i=1, N)])) stop 8
+ if (any (y3 /= [(42, i=1, N)])) stop 9
+
+ call optional_array (y1, y2, y3, is_present=.true.)
+ if (any (y1 /= [(42.24+i, i=1, N)])) stop 10
+ if (any (y2 /= [(42.24+2*i, i=1, N)])) stop 11
+ if (any (y3 /= [(42+3*i, i=1, N)])) stop 12
+
+contains
+ subroutine optional_scalar (a, is_present)
+ integer, optional :: a
+ logical, value :: is_present
+
+ !$omp target data map(a) use_device_addr(a)
+ !$omp target has_device_addr(a)
+ if (is_present) a = 42
+ !$omp end target
+ !$omp end target data
+ end subroutine optional_scalar
+
+ subroutine optional_ptr (a, is_present)
+ integer, pointer, optional :: a
+ logical, value :: is_present
+ !$omp target data map(a) use_device_addr(a)
+ !$omp target has_device_addr(a)
+ if (is_present) a = 84
+ !$omp end target
+ !$omp end target data
+ end subroutine optional_ptr
+
+ subroutine optional_array (a, b, c, is_present)
+ real, optional :: a(:), b(*)
+ integer, optional, pointer, intent(in) :: c(:)
+ logical, value :: is_present
+ integer :: i
+
+ !$omp target data map(a, b(:N), c) use_device_addr(a, b, c)
+ !$omp target has_device_addr(a, b, c)
+ if (is_present) then
+ if (lbound(a,dim=1) /= 1 .or. ubound(a,dim=1) /= N) stop 21
+ if (lbound(b,dim=1) /= 1) stop 22
+ if (lbound(c,dim=1) /= N .or. ubound(c,dim=1) /= 2*N-1) stop 23
+ if (any (a /= [(42.24, i = 1, N)])) stop 24
+ if (any (b(:N) /= [(42.24, i = 1, N)])) stop 25
+ if (any (c /= [(42, i = 1, N)])) stop 26
+ a = [(42.24+i, i=1, N)]
+ b(:N) = [(42.24+2*i, i=1, N)]
+ c = [(42+3*i, i=1, N)]
+ end if
+ !$omp end target
+ !$omp end target data
+ end subroutine optional_array
+
+end program main
--- /dev/null
+! Test allocatables in HAS_DEVICE_ADDR.
+
+program main
+ use omp_lib
+ use iso_c_binding
+ implicit none
+
+ integer, parameter :: N = 5
+ integer, allocatable :: x
+ integer, allocatable :: y(:)
+ call scalar_dummy (x)
+ call array_dummy (y)
+ call array_dummy_optional (y)
+ call array_dummy_optional ()
+
+contains
+ subroutine scalar_dummy (a)
+ integer, allocatable :: a
+
+ allocate (a)
+ a = 24
+
+ !$omp target data map(a) use_device_addr(a)
+ !$omp target has_device_addr(a)
+ a = 42
+ !$omp end target
+ !$omp end target data
+ if (a /= 42) stop 1
+
+ deallocate (a)
+ end subroutine scalar_dummy
+
+ subroutine array_dummy (a)
+ integer, allocatable :: a(:)
+ integer :: i
+
+ allocate (a(N))
+ a = 42
+
+ !$omp target data map(a) use_device_addr(a)
+ !$omp target has_device_addr(a)
+ a = [(i, i=1, N)]
+ !$omp end target
+ !$omp end target data
+ if (any (a /= [(i, i=1, N)])) stop 2
+
+ deallocate (a)
+ end subroutine array_dummy
+
+ subroutine array_dummy_optional (a)
+ integer, optional, allocatable :: a(:)
+ integer :: i
+
+ if (present (a)) then
+ allocate (a(N))
+ a = 42
+ end if
+
+ !$omp target data map(a) use_device_addr(a)
+ !$omp target has_device_addr(a)
+ if (present (a)) a = [(i, i=1, N)]
+ !$omp end target
+ !$omp end target data
+
+ if (present (a)) then
+ if (any (a /= [(i, i=1, N)])) stop 2
+ deallocate (a)
+ end if
+ end subroutine array_dummy_optional
+
+end program main
projects / test_jj.git / commitdiff