diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index ae5094b63337e90591812c7dde43633238544d55..57e9ba14d4a60c67d28b27e25d3fbd14e64f7b8e 100644
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,11 @@
+2016-05-24 Cesar Philippidis <cesar@codesourcery.com>
+
+ * gimplify.c (omp_notice_variable): Use zero-length arrays for data
+ pointers inside OACC_DATA regions.
+ (gimplify_scan_omp_clauses): Prune firstprivate clause associated
+ with OACC_DATA, OACC_ENTER_DATA and OACC_EXIT data regions.
+ (gimplify_adjust_omp_clauses): Fix typo in comment.
+
2016-05-24 Michael Meissner <meissner@linux.vnet.ibm.com>
* config/rs6000/altivec.md (VParity): New mode iterator for vector
diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog
index 3d69cd5794966f6bebad08b1760f47e00b1d2d35..2e3eeb1af3f7592b72ad462db49eec73375221ad 100644
--- a/gcc/c/ChangeLog
+++ b/gcc/c/ChangeLog
@@ -1,3 +1,15 @@
+2016-05-24 Cesar Philippidis <cesar@codesourcery.com>
+
+ * c-parser.c (c_parser_oacc_declare): Add support for
+ GOMP_MAP_FIRSTPRIVATE_POINTER.
+ * c-typeck.c (handle_omp_array_sections_1): Replace bool is_omp
+ argument with enum c_omp_region_type ort.
+ (handle_omp_array_sections): Likewise. Update call to
+ handle_omp_array_sections_1.
+ (c_finish_omp_clauses): Add specific errors and warning messages for
+ OpenACC. Use firsrtprivate pointers for OpenACC subarrays. Update
+ call to handle_omp_array_sections.
+
2016-05-24 Thomas Schwinge <thomas@codesourcery.com>
* c-parser.c (c_parser_oacc_routine): Tighten syntax checks.
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 1bc5eed7f55cb7e5928a3583ca47def96b330cfd..1cf4fb452aa8e66536789aa1c69909a487a50502 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -13602,6 +13602,7 @@ c_parser_oacc_declare (c_parser *parser)
switch (OMP_CLAUSE_MAP_KIND (t))
{
+ case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_FORCE_ALLOC:
case GOMP_MAP_FORCE_TO:
case GOMP_MAP_FORCE_DEVICEPTR:
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index 74bad2acc5d5e95fd39e59f7a14b302e54e3e44e..1520c20c653f42fdc6221ebda436d1d2c61f8ff8 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -11940,7 +11940,7 @@ c_finish_omp_cancellation_point (location_t loc, tree clauses)
static tree
handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
bool &maybe_zero_len, unsigned int &first_non_one,
- bool is_omp)
+ enum c_omp_region_type ort)
{
tree ret, low_bound, length, type;
if (TREE_CODE (t) != TREE_LIST)
@@ -11949,7 +11949,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
return error_mark_node;
ret = t;
if (TREE_CODE (t) == COMPONENT_REF
- && is_omp
+ && ort == C_ORT_OMP
&& (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM))
@@ -11996,7 +11996,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
}
ret = handle_omp_array_sections_1 (c, TREE_CHAIN (t), types,
- maybe_zero_len, first_non_one, is_omp);
+ maybe_zero_len, first_non_one, ort);
if (ret == error_mark_node || ret == NULL_TREE)
return ret;
@@ -12227,14 +12227,14 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
/* Handle array sections for clause C. */
static bool
-handle_omp_array_sections (tree c, bool is_omp)
+handle_omp_array_sections (tree c, enum c_omp_region_type ort)
{
bool maybe_zero_len = false;
unsigned int first_non_one = 0;
auto_vec<tree, 10> types;
tree first = handle_omp_array_sections_1 (c, OMP_CLAUSE_DECL (c), types,
maybe_zero_len, first_non_one,
- is_omp);
+ ort);
if (first == error_mark_node)
return true;
if (first == NULL_TREE)
@@ -12427,7 +12427,7 @@ handle_omp_array_sections (tree c, bool is_omp)
&& TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE))
return false;
gcc_assert (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FORCE_DEVICEPTR);
- if (is_omp)
+ if (ort == C_ORT_OMP || ort == C_ORT_ACC)
switch (OMP_CLAUSE_MAP_KIND (c))
{
case GOMP_MAP_ALLOC:
@@ -12445,7 +12445,7 @@ handle_omp_array_sections (tree c, bool is_omp)
break;
}
tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
- if (!is_omp)
+ 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)
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
@@ -12520,7 +12520,7 @@ tree
c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
{
bitmap_head generic_head, firstprivate_head, lastprivate_head;
- bitmap_head aligned_head, map_head, map_field_head;
+ bitmap_head aligned_head, map_head, map_field_head, oacc_reduction_head;
tree c, t, type, *pc;
tree simdlen = NULL_TREE, safelen = NULL_TREE;
bool branch_seen = false;
@@ -12537,6 +12537,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
bitmap_initialize (&aligned_head, &bitmap_default_obstack);
bitmap_initialize (&map_head, &bitmap_default_obstack);
bitmap_initialize (&map_field_head, &bitmap_default_obstack);
+ bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack);
for (pc = &clauses, c = clauses; c ; c = *pc)
{
@@ -12560,7 +12561,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
t = OMP_CLAUSE_DECL (c);
if (TREE_CODE (t) == TREE_LIST)
{
- if (handle_omp_array_sections (c, ort & C_ORT_OMP))
+ if (handle_omp_array_sections (c, ort))
{
remove = true;
break;
@@ -12874,6 +12875,17 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
+ else if (ort == C_ORT_ACC
+ && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
+ {
+ if (bitmap_bit_p (&oacc_reduction_head, DECL_UID (t)))
+ {
+ error ("%qD appears more than once in reduction clauses", t);
+ remove = true;
+ }
+ else
+ bitmap_set_bit (&oacc_reduction_head, DECL_UID (t));
+ }
else if (bitmap_bit_p (&generic_head, DECL_UID (t))
|| bitmap_bit_p (&firstprivate_head, DECL_UID (t))
|| bitmap_bit_p (&lastprivate_head, DECL_UID (t)))
@@ -12885,7 +12897,10 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_PRIVATE
&& bitmap_bit_p (&map_head, DECL_UID (t)))
{
- error ("%qD appears both in data and map clauses", t);
+ if (ort == C_ORT_ACC)
+ error ("%qD appears more than once in data clauses", t);
+ else
+ error ("%qD appears both in data and map clauses", t);
remove = true;
}
else
@@ -12911,7 +12926,10 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
}
else if (bitmap_bit_p (&map_head, DECL_UID (t)))
{
- error ("%qD appears both in data and map clauses", t);
+ if (ort == C_ORT_ACC)
+ error ("%qD appears more than once in data clauses", t);
+ else
+ error ("%qD appears both in data and map clauses", t);
remove = true;
}
else
@@ -13004,7 +13022,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
}
if (TREE_CODE (t) == TREE_LIST)
{
- if (handle_omp_array_sections (c, ort & C_ORT_OMP))
+ if (handle_omp_array_sections (c, ort))
remove = true;
break;
}
@@ -13027,7 +13045,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
t = OMP_CLAUSE_DECL (c);
if (TREE_CODE (t) == TREE_LIST)
{
- if (handle_omp_array_sections (c, ort & C_ORT_OMP))
+ if (handle_omp_array_sections (c, ort))
remove = true;
else
{
@@ -13054,6 +13072,9 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
error ("%qD appears more than once in motion"
" clauses", t);
+ else if (ort == C_ORT_ACC)
+ error ("%qD appears more than once in data"
+ " clauses", t);
else
error ("%qD appears more than once in map"
" clauses", t);
@@ -13155,7 +13176,10 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
}
else if (bitmap_bit_p (&map_head, DECL_UID (t)))
{
- error ("%qD appears both in data and map clauses", t);
+ if (ort == C_ORT_ACC)
+ error ("%qD appears more than once in data clauses", t);
+ else
+ error ("%qD appears both in data and map clauses", t);
remove = true;
}
else
@@ -13165,6 +13189,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
{
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
error ("%qD appears more than once in motion clauses", t);
+ else if (ort == C_ORT_ACC)
+ error ("%qD appears more than once in data clauses", t);
else
error ("%qD appears more than once in map clauses", t);
remove = true;
@@ -13172,7 +13198,10 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
else if (bitmap_bit_p (&generic_head, DECL_UID (t))
|| bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
{
- error ("%qD appears both in data and map clauses", t);
+ if (ort == C_ORT_ACC)
+ error ("%qD appears more than once in data clauses", t);
+ else
+ error ("%qD appears both in data and map clauses", t);
remove = true;
}
else
diff --git a/gcc/cp/ChangeLog b/gcc/cp/ChangeLog
index 8010cb70b94fcf23d2c347a306cf23262c4371c5..9f618a4f2ffc417760036d6aa3ddd0d3e7155224 100644
--- a/gcc/cp/ChangeLog
+++ b/gcc/cp/ChangeLog
@@ -1,3 +1,16 @@
+2016-05-24 Cesar Philippidis <cesar@codesourcery.com>
+
+ * parser.c (cp_parser_oacc_declare): Add support for
+ GOMP_MAP_FIRSTPRIVATE_POINTER.
+ * semantics.c (handle_omp_array_sections_1): Replace bool is_omp
+ argument with enum c_omp_region_type ort. Don't privatize OpenACC
+ non-static members.
+ (handle_omp_array_sections): Replace bool is_omp argument with enum
+ c_omp_region_type ort. Update call to handle_omp_array_sections_1.
+ (finish_omp_clauses): Add specific errors and warning messages for
+ OpenACC. Use firsrtprivate pointers for OpenACC subarrays. Update
+ call to handle_omp_array_sections.
+
2016-05-24 Jason Merrill <jason@redhat.com>
PR c++/70584
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index a68a510c6601519633d3f8a3ff0bcce655b153b6..d21230fdd8d6a0bbf9ec6048d3be9d55d9d4928c 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -35214,6 +35214,7 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
gcc_assert (OMP_CLAUSE_CODE (t) == OMP_CLAUSE_MAP);
switch (OMP_CLAUSE_MAP_KIND (t))
{
+ case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_FORCE_ALLOC:
case GOMP_MAP_FORCE_TO:
case GOMP_MAP_FORCE_DEVICEPTR:
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 21ef42f5534da50f66958c6483f9140cdc0b8b83..8a3e7fd619b33ef88241cb390b3a29d387994b6c 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -4472,7 +4472,7 @@ omp_privatize_field (tree t, bool shared)
static tree
handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
bool &maybe_zero_len, unsigned int &first_non_one,
- bool is_omp)
+ enum c_omp_region_type ort)
{
tree ret, low_bound, length, type;
if (TREE_CODE (t) != TREE_LIST)
@@ -4484,7 +4484,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
t = TREE_OPERAND (t, 0);
ret = t;
if (TREE_CODE (t) == COMPONENT_REF
- && is_omp
+ && ort == C_ORT_OMP
&& (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM)
@@ -4545,11 +4545,12 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
return ret;
}
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION
+ if (ort == C_ORT_OMP
+ && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION
&& TREE_CODE (TREE_CHAIN (t)) == FIELD_DECL)
TREE_CHAIN (t) = omp_privatize_field (TREE_CHAIN (t), false);
ret = handle_omp_array_sections_1 (c, TREE_CHAIN (t), types,
- maybe_zero_len, first_non_one, is_omp);
+ maybe_zero_len, first_non_one, ort);
if (ret == error_mark_node || ret == NULL_TREE)
return ret;
@@ -4792,14 +4793,14 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
/* Handle array sections for clause C. */
static bool
-handle_omp_array_sections (tree c, bool is_omp)
+handle_omp_array_sections (tree c, enum c_omp_region_type ort)
{
bool maybe_zero_len = false;
unsigned int first_non_one = 0;
auto_vec<tree, 10> types;
tree first = handle_omp_array_sections_1 (c, OMP_CLAUSE_DECL (c), types,
maybe_zero_len, first_non_one,
- is_omp);
+ ort);
if (first == error_mark_node)
return true;
if (first == NULL_TREE)
@@ -4988,7 +4989,7 @@ handle_omp_array_sections (tree c, bool is_omp)
|| (TREE_CODE (t) == COMPONENT_REF
&& TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE))
return false;
- if (is_omp)
+ if (ort == C_ORT_OMP || ort == C_ORT_ACC)
switch (OMP_CLAUSE_MAP_KIND (c))
{
case GOMP_MAP_ALLOC:
@@ -5007,7 +5008,7 @@ handle_omp_array_sections (tree c, bool is_omp)
}
tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
OMP_CLAUSE_MAP);
- if (!is_omp)
+ 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)
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
@@ -5774,7 +5775,7 @@ tree
finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
{
bitmap_head generic_head, firstprivate_head, lastprivate_head;
- bitmap_head aligned_head, map_head, map_field_head;
+ bitmap_head aligned_head, map_head, map_field_head, oacc_reduction_head;
tree c, t, *pc;
tree safelen = NULL_TREE;
bool branch_seen = false;
@@ -5788,6 +5789,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
bitmap_initialize (&aligned_head, &bitmap_default_obstack);
bitmap_initialize (&map_head, &bitmap_default_obstack);
bitmap_initialize (&map_field_head, &bitmap_default_obstack);
+ bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack);
for (pc = &clauses, c = clauses; c ; c = *pc)
{
@@ -5807,8 +5809,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
t = OMP_CLAUSE_DECL (c);
if (TREE_CODE (t) == TREE_LIST)
{
- if (handle_omp_array_sections (c, ((ort & C_ORT_OMP_DECLARE_SIMD)
- == C_ORT_OMP)))
+ if (handle_omp_array_sections (c, ort))
{
remove = true;
break;
@@ -6018,6 +6019,17 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
+ else if (ort == C_ORT_ACC
+ && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
+ {
+ if (bitmap_bit_p (&oacc_reduction_head, DECL_UID (t)))
+ {
+ error ("%qD appears more than once in reduction clauses", t);
+ remove = true;
+ }
+ else
+ bitmap_set_bit (&oacc_reduction_head, DECL_UID (t));
+ }
else if (bitmap_bit_p (&generic_head, DECL_UID (t))
|| bitmap_bit_p (&firstprivate_head, DECL_UID (t))
|| bitmap_bit_p (&lastprivate_head, DECL_UID (t)))
@@ -6028,7 +6040,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_PRIVATE
&& bitmap_bit_p (&map_head, DECL_UID (t)))
{
- error ("%qD appears both in data and map clauses", t);
+ if (ort == C_ORT_ACC)
+ error ("%qD appears more than once in data clauses", t);
+ else
+ error ("%qD appears both in data and map clauses", t);
remove = true;
}
else
@@ -6038,7 +6053,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
handle_field_decl:
if (!remove
&& TREE_CODE (t) == FIELD_DECL
- && t == OMP_CLAUSE_DECL (c))
+ && t == OMP_CLAUSE_DECL (c)
+ && ort != C_ORT_ACC)
{
OMP_CLAUSE_DECL (c)
= omp_privatize_field (t, (OMP_CLAUSE_CODE (c)
@@ -6054,7 +6070,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
omp_note_field_privatization (t, OMP_CLAUSE_DECL (c));
else
t = OMP_CLAUSE_DECL (c);
- if (t == current_class_ptr)
+ if (ort != C_ORT_ACC && t == current_class_ptr)
{
error ("%<this%> allowed in OpenMP only in %<declare simd%>"
" clauses");
@@ -6081,7 +6097,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
}
else if (bitmap_bit_p (&map_head, DECL_UID (t)))
{
- error ("%qD appears both in data and map clauses", t);
+ if (ort == C_ORT_ACC)
+ error ("%qD appears more than once in data clauses", t);
+ else
+ error ("%qD appears both in data and map clauses", t);
remove = true;
}
else
@@ -6529,8 +6548,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
}
if (TREE_CODE (t) == TREE_LIST)
{
- if (handle_omp_array_sections (c, ((ort & C_ORT_OMP_DECLARE_SIMD)
- == C_ORT_OMP)))
+ if (handle_omp_array_sections (c, ort))
remove = true;
break;
}
@@ -6564,8 +6582,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
t = OMP_CLAUSE_DECL (c);
if (TREE_CODE (t) == TREE_LIST)
{
- if (handle_omp_array_sections (c, ((ort & C_ORT_OMP_DECLARE_SIMD)
- == C_ORT_OMP)))
+ if (handle_omp_array_sections (c, ort))
remove = true;
else
{
@@ -6594,6 +6611,9 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
error ("%qD appears more than once in motion"
" clauses", t);
+ else if (ort == C_ORT_ACC)
+ error ("%qD appears more than once in data"
+ " clauses", t);
else
error ("%qD appears more than once in map"
" clauses", t);
@@ -6681,7 +6701,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
- else if (t == current_class_ptr)
+ else if (ort != C_ORT_ACC && t == current_class_ptr)
{
error ("%<this%> allowed in OpenMP only in %<declare simd%>"
" clauses");
@@ -6730,7 +6750,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
}
else if (bitmap_bit_p (&map_head, DECL_UID (t)))
{
- error ("%qD appears both in data and map clauses", t);
+ if (ort == C_ORT_ACC)
+ error ("%qD appears more than once in data clauses", t);
+ else
+ error ("%qD appears both in data and map clauses", t);
remove = true;
}
else
@@ -6740,6 +6763,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
{
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
error ("%qD appears more than once in motion clauses", t);
+ if (ort == C_ORT_ACC)
+ error ("%qD appears more than once in data clauses", t);
else
error ("%qD appears more than once in map clauses", t);
remove = true;
@@ -6747,7 +6772,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
else if (bitmap_bit_p (&generic_head, DECL_UID (t))
|| bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
{
- error ("%qD appears both in data and map clauses", t);
+ if (ort == C_ORT_ACC)
+ error ("%qD appears more than once in data clauses", t);
+ else
+ error ("%qD appears both in data and map clauses", t);
remove = true;
}
else
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 67394e36407c29e237ac0225ba181879cbea5479..8316bb8881f22745b705d6bb9f2e3a1af80d43f4 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -6280,6 +6280,9 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
error ("variable %qE declared in enclosing "
"%<host_data%> region", DECL_NAME (decl));
nflags |= GOVD_MAP;
+ if (octx->region_type == ORT_ACC_DATA
+ && (n2->value & GOVD_MAP_0LEN_ARRAY))
+ nflags |= GOVD_MAP_0LEN_ARRAY;
goto found_outer;
}
}
@@ -6855,9 +6858,14 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
{
case OMP_TARGET:
break;
+ case OACC_DATA:
+ if (TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE)
+ break;
case OMP_TARGET_DATA:
case OMP_TARGET_ENTER_DATA:
case OMP_TARGET_EXIT_DATA:
+ case OACC_ENTER_DATA:
+ case OACC_EXIT_DATA:
case OACC_HOST_DATA:
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
|| (OMP_CLAUSE_MAP_KIND (c)
@@ -7311,6 +7319,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
omp_notice_variable (outer_ctx, t, true);
}
}
+ if (code == OACC_DATA
+ && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ flags |= GOVD_MAP_0LEN_ARRAY;
omp_add_variable (ctx, decl, flags);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION
&& OMP_CLAUSE_REDUCTION_PLACEHOLDER (c))
@@ -7569,6 +7581,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
gcc_unreachable ();
}
+ if (code == OACC_DATA
+ && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ remove = true;
if (remove)
*list_p = OMP_CLAUSE_CHAIN (c);
else
@@ -8029,7 +8045,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
break;
}
decl = OMP_CLAUSE_DECL (c);
- /* Data clasues associated with acc parallel reductions must be
+ /* Data clauses associated with acc parallel reductions must be
compatible with present_or_copy. Warn and adjust the clause
if that is not the case. */
if (ctx->region_type == ORT_ACC_PARALLEL)
diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index 1cde48184c04254053fe06bd6d241d972278a4ee..bb082a177ced45d05eac794609250dfc2dbb319a 100644
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,3 +1,21 @@
+2016-05-24 Cesar Philippidis <cesar@codesourcery.com>
+
+ * c-c++-common/goacc/data-clause-duplicate-1.c: Adjust test.
+ * c-c++-common/goacc/deviceptr-1.c: Likewise.
+ * c-c++-common/goacc/kernels-alias-3.c: Likewise.
+ * c-c++-common/goacc/kernels-alias-4.c: Likewise.
+ * c-c++-common/goacc/kernels-alias-5.c: Likewise.
+ * c-c++-common/goacc/kernels-alias-8.c: Likewise.
+ * c-c++-common/goacc/kernels-alias-ipa-pta-3.c: Likewise.
+ * c-c++-common/goacc/pcopy.c: Likewise.
+ * c-c++-common/goacc/pcopyin.c: Likewise.
+ * c-c++-common/goacc/pcopyout.c: Likewise.
+ * c-c++-common/goacc/pcreate.c: Likewise.
+ * c-c++-common/goacc/pr70688.c: New test.
+ * c-c++-common/goacc/present-1.c: Adjust test.
+ * c-c++-common/goacc/reduction-5.c: Likewise.
+ * g++.dg/goacc/data-1.C: New test.
+
2016-05-24 Michael Meissner <meissner@linux.vnet.ibm.com>
* gcc.target/powerpc/p9-vparity.c: New file to check ISA 3.0
diff --git a/gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c b/gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c
index 7a1cf683f244706db48181a213bd7f593a25e893..6245beb20794e2c60f437b05388b7c9f0fc44504 100644
--- a/gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c
@@ -2,12 +2,12 @@ void
fun (void)
{
float *fp;
-#pragma acc parallel copy(fp[0:2],fp[0:2]) /* { dg-error "'fp' appears more than once in map clauses" } */
+#pragma acc parallel copy(fp[0:2],fp[0:2]) /* { dg-error "'fp' appears more than once in data clauses" } */
;
-#pragma acc kernels present_or_copyin(fp[3]) present_or_copyout(fp[7:4]) /* { dg-error "'fp' appears more than once in map clauses" } */
+#pragma acc kernels present_or_copyin(fp[3]) present_or_copyout(fp[7:4]) /* { dg-error "'fp' appears more than once in data clauses" } */
;
-#pragma acc data create(fp[:10]) deviceptr(fp) /* { dg-error "'fp' appears more than once in map clauses" } */
+#pragma acc data create(fp[:10]) deviceptr(fp) /* { dg-error "'fp' appears more than once in data clauses" } */
;
-#pragma acc data create(fp) present(fp) /* { dg-error "'fp' appears more than once in map clauses" } */
+#pragma acc data create(fp) present(fp) /* { dg-error "'fp' appears more than once in data clauses" } */
;
}
diff --git a/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c b/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
index 08ddb1072ff8e4b6c8c966cf5682309e16964928..3aa0e8a115364d897f82c4570e2e1d9e65d39ac4 100644
--- a/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
@@ -47,7 +47,7 @@ fun2 (void)
/* { dg-error "'u' undeclared" "u undeclared" { target *-*-* } 46 } */
/* { dg-error "'fun2' is not a variable" "fun2 not a variable" { target *-*-* } 46 } */
/* { dg-error "'i' is not a pointer variable" "i not a pointer variable" { target *-*-* } 46 } */
- /* { dg-error "'fp' appears more than once in map clauses" "fp more than once" { target *-*-* } 46 } */
+ /* { dg-error "'fp' appears more than once in data clauses" "fp more than once" { target *-*-* } 46 } */
;
}
@@ -55,11 +55,11 @@ void
fun3 (void)
{
float *fp;
-#pragma acc data deviceptr(fp,fp) /* { dg-error "'fp' appears more than once in map clauses" } */
+#pragma acc data deviceptr(fp,fp) /* { dg-error "'fp' appears more than once in data clauses" } */
;
-#pragma acc parallel deviceptr(fp) deviceptr(fp) /* { dg-error "'fp' appears more than once in map clauses" } */
+#pragma acc parallel deviceptr(fp) deviceptr(fp) /* { dg-error "'fp' appears more than once in data clauses" } */
;
-#pragma acc kernels copy(fp) deviceptr(fp) /* { dg-error "'fp' appears more than once in map clauses" } */
+#pragma acc kernels copy(fp) deviceptr(fp) /* { dg-error "'fp' appears more than once in data clauses" } */
;
}
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-3.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-3.c
index 6989c1c1b6d1ac79afb967158848512d9a515b36..2934f124ef69edca1011d289394511a065fa581f 100644
--- a/gcc/testsuite/c-c++-common/goacc/kernels-alias-3.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-3.c
@@ -17,5 +17,5 @@ foo (void)
/* Only the omp_data_i related loads should be annotated with
non-base 0 cliques. */
/* { dg-final { scan-tree-dump-times "clique 1 base 1" 2 "ealias" } } */
-/* { dg-final { scan-tree-dump-times "(?n)clique 1 base 0" 3 "ealias" } } */
+/* { dg-final { scan-tree-dump-times "(?n)clique 1 base 0" 2 "ealias" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-4.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-4.c
index d41802c00db76e43c90d98748f3de903a82cc662..f6ee5b58c7126f6cb86565bf5cab0f3f3c4275e4 100644
--- a/gcc/testsuite/c-c++-common/goacc/kernels-alias-4.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-4.c
@@ -19,5 +19,5 @@ foo (void)
/* Only the omp_data_i related loads should be annotated with
non-base 0 cliques. */
/* { dg-final { scan-tree-dump-times "clique 1 base 1" 2 "ealias" } } */
-/* { dg-final { scan-tree-dump-times "(?n)clique 1 base 0" 3 "ealias" } } */
+/* { dg-final { scan-tree-dump-times "(?n)clique 1 base 0" 2 "ealias" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-5.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-5.c
index 6fefe183203687d31bce5ec8acb93bb0f2ff19dd..74425fb02bc9b0d5b4e2509fa24b331c4e61aa54 100644
--- a/gcc/testsuite/c-c++-common/goacc/kernels-alias-5.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-5.c
@@ -15,5 +15,5 @@ foo (int *a)
/* Only the omp_data_i related loads should be annotated with cliques. */
/* { dg-final { scan-tree-dump-times "clique 1 base 1" 2 "ealias" } } */
-/* { dg-final { scan-tree-dump-times "(?n)clique 1 base 0" 4 "ealias" } } */
+/* { dg-final { scan-tree-dump-times "(?n)clique 1 base 0" 2 "ealias" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-8.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-8.c
index 3b91acd27f57d419d2d017eeeaa7fa27f6809414..69200ccf192556cb762f10a668a326368581e4bf 100644
--- a/gcc/testsuite/c-c++-common/goacc/kernels-alias-8.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-8.c
@@ -7,7 +7,7 @@ extern void *acc_copyin (void *, size_t);
void
foo (int *a, size_t n)
{
- int *p = (int *)acc_copyin (&a, n);
+ int *p = (int *)acc_copyin (a, n);
#pragma acc kernels deviceptr (p) pcopy(a[0:n])
{
@@ -18,5 +18,5 @@ foo (int *a, size_t n)
/* Only the omp_data_i related loads should be annotated with cliques. */
/* { dg-final { scan-tree-dump-times "clique 1 base 1" 2 "ealias" } } */
-/* { dg-final { scan-tree-dump-times "(?n)clique 1 base 0" 3 "ealias" } } */
+/* { dg-final { scan-tree-dump-times "(?n)clique 1 base 0" 2 "ealias" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-3.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-3.c
index 1eb56eb9e62d46e72268a33a4a41b1a2ce892062..1ea0e7303d900a0d8935255fe7fb6037716372ab 100644
--- a/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-3.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-3.c
@@ -31,6 +31,5 @@ foo (void)
free (c);
}
-/* { dg-final { scan-tree-dump-times "(?n)= 0;$" 1 "optimized" } } */
-/* { dg-final { scan-tree-dump-times "(?n)= 1;$" 1 "optimized" } } */
-/* { dg-final { scan-tree-dump-times "(?n)= \\*a" 1 "optimized" } } */
+/* { dg-final { scan-tree-dump-times "(?n)= 0;$" 1 "optimized" { target c } } } */
+/* { dg-final { scan-tree-dump-times "(?n)= 1;$" 1 "optimized" { target c } } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/pcopy.c b/gcc/testsuite/c-c++-common/goacc/pcopy.c
index 02c438353e78a10bbec249b43472b6111c0b2e64..0e0aad55a6806d96c19212f223c6fcbd15211ecd 100644
--- a/gcc/testsuite/c-c++-common/goacc/pcopy.c
+++ b/gcc/testsuite/c-c++-common/goacc/pcopy.c
@@ -7,4 +7,4 @@ f (char *cp)
;
}
-/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(tofrom:\\*\\(cp \\+ 3\\) \\\[len: 5]\\) map\\(alloc:cp \\\[pointer assign, bias: 3]\\)" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(tofrom:\\*\\(cp \\+ 3\\) \\\[len: 5]\\) map\\(firstprivate:cp \\\[pointer assign, bias: 3]\\)" 1 "original" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/pcopyin.c b/gcc/testsuite/c-c++-common/goacc/pcopyin.c
index 10911fcbb4d33a4de3c88e01fad39b57df5195bf..3085251deb653cb3555b4e849065ef9b6ef1e592 100644
--- a/gcc/testsuite/c-c++-common/goacc/pcopyin.c
+++ b/gcc/testsuite/c-c++-common/goacc/pcopyin.c
@@ -7,4 +7,4 @@ f (char *cp)
;
}
-/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(to:\\*\\(cp \\+ 4\\) \\\[len: 6]\\) map\\(alloc:cp \\\[pointer assign, bias: 4]\\)" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(to:\\*\\(cp \\+ 4\\) \\\[len: 6]\\) map\\(firstprivate:cp \\\[pointer assign, bias: 4]\\)" 1 "original" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/pcopyout.c b/gcc/testsuite/c-c++-common/goacc/pcopyout.c
index 703ac2f9010929ea33f8ab4f2eaac7e0794565cc..47c454ce92a5eca82fed0ad62dc8cc0ac91a01a4 100644
--- a/gcc/testsuite/c-c++-common/goacc/pcopyout.c
+++ b/gcc/testsuite/c-c++-common/goacc/pcopyout.c
@@ -7,4 +7,4 @@ f (char *cp)
;
}
-/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(from:\\*\\(cp \\+ 5\\) \\\[len: 7]\\) map\\(alloc:cp \\\[pointer assign, bias: 5]\\)" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(from:\\*\\(cp \\+ 5\\) \\\[len: 7]\\) map\\(firstprivate:cp \\\[pointer assign, bias: 5]\\)" 1 "original" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/pcreate.c b/gcc/testsuite/c-c++-common/goacc/pcreate.c
index 00bf1551a5cb31163656c1fcd4327678d7f77620..a403e5a64ffe3ae3c5a2eb6c06c59ac335cb5126 100644
--- a/gcc/testsuite/c-c++-common/goacc/pcreate.c
+++ b/gcc/testsuite/c-c++-common/goacc/pcreate.c
@@ -7,4 +7,4 @@ f (char *cp)
;
}
-/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(alloc:\\*\\(cp \\+ 6\\) \\\[len: 8]\\) map\\(alloc:cp \\\[pointer assign, bias: 6]\\)" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(alloc:\\*\\(cp \\+ 6\\) \\\[len: 8]\\) map\\(firstprivate:cp \\\[pointer assign, bias: 6]\\)" 1 "original" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/pr70688.c b/gcc/testsuite/c-c++-common/goacc/pr70688.c
new file mode 100644
index 0000000000000000000000000000000000000000..5a2366540692b491f589afdbb22f8e09f5efc5e2
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/pr70688.c
@@ -0,0 +1,48 @@
+const int n = 100;
+
+int
+private_reduction ()
+{
+ int i, r;
+
+ #pragma acc parallel
+ #pragma acc loop private (r) reduction (+:r)
+ for (i = 0; i < 100; i++)
+ r += 10;
+
+ return r;
+}
+
+int
+parallel_reduction ()
+{
+ int sum = 0;
+ int dummy = 0;
+
+#pragma acc data copy (dummy)
+ {
+#pragma acc parallel num_gangs (10) copy (sum) reduction (+:sum)
+ {
+ int v = 5;
+ sum += 10 + v;
+ }
+ }
+
+ return sum;
+}
+
+int
+main ()
+{
+ int i, s = 0;
+
+#pragma acc parallel num_gangs (10) copy (s) reduction (+:s)
+ for (i = 0; i < n; i++)
+ s += i+1;
+
+#pragma acc parallel num_gangs (10) reduction (+:s) copy (s)
+ for (i = 0; i < n; i++)
+ s += i+1;
+
+ return 0;
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/present-1.c b/gcc/testsuite/c-c++-common/goacc/present-1.c
index 753794852d731b189c58c0574cc79857a87708c5..51362b2f2f16eb98bb940dd340d18ddb30ac94c5 100644
--- a/gcc/testsuite/c-c++-common/goacc/present-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/present-1.c
@@ -7,4 +7,4 @@ f (char *cp)
;
}
-/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(force_present:\\*\\(cp \\+ 7\\) \\\[len: 9]\\) map\\(alloc:cp \\\[pointer assign, bias: 7]\\)" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(force_present:\\*\\(cp \\+ 7\\) \\\[len: 9]\\) map\\(firstprivate:cp \\\[pointer assign, bias: 7]\\)" 1 "original" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/reduction-5.c b/gcc/testsuite/c-c++-common/goacc/reduction-5.c
index 74daad35a5d7a786d1f83f52b899a82590780967..dfdbab9139375a418ca1281c7da8352a528f384e 100644
--- a/gcc/testsuite/c-c++-common/goacc/reduction-5.c
+++ b/gcc/testsuite/c-c++-common/goacc/reduction-5.c
@@ -7,9 +7,9 @@ main(void)
{
int v1;
-#pragma acc parallel reduction(+:v1) private(v1) /* { dg-error "appears more than once in data clauses" } */
+#pragma acc parallel reduction(+:v1) private(v1) /* { dg-error "invalid private reduction" } */
;
-#pragma acc parallel reduction(+:v1) firstprivate(v1) /* { dg-error "appears more than once in data clauses" } */
+#pragma acc parallel reduction(+:v1) firstprivate(v1) /* { dg-error "invalid private reduction" } */
;
return 0;
diff --git a/gcc/testsuite/g++.dg/goacc/data-1.C b/gcc/testsuite/g++.dg/goacc/data-1.C
new file mode 100644
index 0000000000000000000000000000000000000000..54676dc2dffbb477c8e744fc6e1f59b1adc4c32d
--- /dev/null
+++ b/gcc/testsuite/g++.dg/goacc/data-1.C
@@ -0,0 +1,39 @@
+void
+foo (int &a, int (&b)[100], int &n)
+{
+#pragma acc enter data copyin (a, b) async wait
+#pragma acc enter data create (b[20:30]) async wait
+#pragma acc enter data (a) /* { dg-error "expected '#pragma acc' clause before '\\\(' token" } */
+#pragma acc enter data create (b(1:10)) /* { dg-error "expected '\\\)' before '\\\(' token" } */
+#pragma acc exit data delete (a) if (0)
+#pragma acc exit data copyout (b) if (a)
+#pragma acc exit data delete (b)
+#pragma acc enter /* { dg-error "expected 'data' in" } */
+#pragma acc exit /* { dg-error "expected 'data' in" } */
+#pragma acc enter data /* { dg-error "has no data movement clause" } */
+#pragma acc exit data /* { dg-error "has no data movement clause" } */
+#pragma acc enter Data /* { dg-error "invalid pragma before" } */
+#pragma acc exit copyout (b) /* { dg-error "invalid pragma before" } */
+}
+
+template<typename T>
+void
+foo (T &a, T (&b)[100], T &n)
+{
+#pragma acc enter data copyin (a, b) async wait
+#pragma acc enter data create (b[20:30]) async wait
+#pragma acc enter data (a) /* { dg-error "expected '#pragma acc' clause before '\\\(' token" } */
+#pragma acc enter data create (b(1:10)) /* { dg-error "expected '\\\)' before '\\\(' token" } */
+#pragma acc exit data delete (a) if (0)
+#pragma acc exit data copyout (b) if (a)
+#pragma acc exit data delete (b)
+#pragma acc enter /* { dg-error "expected 'data' in" } */
+#pragma acc exit /* { dg-error "expected 'data' in" } */
+#pragma acc enter data /* { dg-error "has no data movement clause" } */
+#pragma acc exit data /* { dg-error "has no data movement clause" } */
+#pragma acc enter Data /* { dg-error "invalid pragma before" } */
+#pragma acc exit copyout (b) /* { dg-error "invalid pragma before" } */
+}
+
+/* { dg-error "has no data movement clause" "" { target *-*-* } 6 } */
+/* { dg-error "has no data movement clause" "" { target *-*-* } 25 } */
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 6a390879ff160713f6057019d9355ed2b9326fc8..7ad7ff49ff85687e23408b9c76e86f472588d868 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,48 @@
+2016-05-24 Cesar Philippidis <cesar@codesourcery.com>
+
+ * oacc-mem.c (acc_malloc): Update handling of shared-memory targets.
+ (acc_free): Likewise.
+ (acc_memcpy_to_device): Likewise.
+ (acc_memcpy_from_device): Likewise.
+ (acc_deviceptr): Likewise.
+ (acc_hostptr): Likewise.
+ (acc_is_present): Likewise.
+ (acc_map_data): Likewise.
+ (acc_unmap_data): Likewise.
+ (present_create_copy): Likewise.
+ (delete_copyout): Likewise.
+ (update_dev_host): Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c: Remove xfail.
+ * testsuite/libgomp.oacc-c-c++-common/data-2-lib.c: New test.
+ * testsuite/libgomp.oacc-c-c++-common/data-2.c: Adjust test.
+ * testsuite/libgomp.oacc-c-c++-common/data-3.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/enter_exit-lib.c: New test.
+ * testsuite/libgomp.oacc-c-c++-common/lib-13.c: Adjust test so that
+ it only runs on nvptx targets.
+ * testsuite/libgomp.oacc-c-c++-common/lib-14.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-15.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-16.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-17.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-18.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-20.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-21.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-22.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-23.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-24.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-25.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-28.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-29.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-30.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-34.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-42.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-43.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-44.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-47.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-48.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-52.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-53.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/lib-54.c: Likewise.
+
2016-05-23 Martin Jambor <mjambor@suse.cz>
* testsuite/libgomp.hsa.c/switch-sbr-2.c: New test.
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index ce1905cb271dbd2347f0e950400c70fd7f6e83a4..665e208cd36e1178fbcfc51ba7f42cc5a7b9e7d7 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -32,6 +32,7 @@
#include "gomp-constants.h"
#include "oacc-int.h"
#include <stdint.h>
+#include <string.h>
#include <assert.h>
/* Return block containing [H->S), or NULL if not contained. The device lock
@@ -104,6 +105,9 @@ acc_malloc (size_t s)
assert (thr->dev);
+ if (thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
+ return malloc (s);
+
return thr->dev->alloc_func (thr->dev->target_id, s);
}
@@ -124,6 +128,9 @@ acc_free (void *d)
struct gomp_device_descr *acc_dev = thr->dev;
+ if (acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
+ return free (d);
+
gomp_mutex_lock (&acc_dev->lock);
/* We don't have to call lazy open here, as the ptr value must have
@@ -154,6 +161,12 @@ acc_memcpy_to_device (void *d, void *h, size_t s)
assert (thr && thr->dev);
+ if (thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
+ {
+ memmove (d, h, s);
+ return;
+ }
+
thr->dev->host2dev_func (thr->dev->target_id, d, h, s);
}
@@ -166,6 +179,12 @@ acc_memcpy_from_device (void *h, void *d, size_t s)
assert (thr && thr->dev);
+ if (thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
+ {
+ memmove (h, d, s);
+ return;
+ }
+
thr->dev->dev2host_func (thr->dev->target_id, h, d, s);
}
@@ -184,6 +203,9 @@ acc_deviceptr (void *h)
struct goacc_thread *thr = goacc_thread ();
struct gomp_device_descr *dev = thr->dev;
+ if (thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
+ return h;
+
gomp_mutex_lock (&dev->lock);
n = lookup_host (dev, h, 1);
@@ -218,6 +240,9 @@ acc_hostptr (void *d)
struct goacc_thread *thr = goacc_thread ();
struct gomp_device_descr *acc_dev = thr->dev;
+ if (thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
+ return d;
+
gomp_mutex_lock (&acc_dev->lock);
n = lookup_dev (acc_dev->openacc.data_environ, d, 1);
@@ -252,6 +277,9 @@ acc_is_present (void *h, size_t s)
struct goacc_thread *thr = goacc_thread ();
struct gomp_device_descr *acc_dev = thr->dev;
+ if (thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
+ return h != NULL;
+
gomp_mutex_lock (&acc_dev->lock);
n = lookup_host (acc_dev, h, s);
@@ -271,7 +299,7 @@ acc_is_present (void *h, size_t s)
void
acc_map_data (void *h, void *d, size_t s)
{
- struct target_mem_desc *tgt;
+ struct target_mem_desc *tgt = NULL;
size_t mapnum = 1;
void *hostaddrs = h;
void *devaddrs = d;
@@ -287,9 +315,6 @@ acc_map_data (void *h, void *d, size_t s)
{
if (d != h)
gomp_fatal ("cannot map data on shared-memory system");
-
- tgt = gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, true,
- GOMP_MAP_VARS_OPENACC);
}
else
{
@@ -335,6 +360,10 @@ acc_unmap_data (void *h)
/* No need to call lazy open, as the address must have been mapped. */
+ /* This is a no-op on shared-memory targets. */
+ if (acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
+ return;
+
size_t host_size;
gomp_mutex_lock (&acc_dev->lock);
@@ -405,6 +434,9 @@ present_create_copy (unsigned f, void *h, size_t s)
struct goacc_thread *thr = goacc_thread ();
struct gomp_device_descr *acc_dev = thr->dev;
+ if (acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
+ return h;
+
gomp_mutex_lock (&acc_dev->lock);
n = lookup_host (acc_dev, h, s);
@@ -496,6 +528,9 @@ delete_copyout (unsigned f, void *h, size_t s)
struct goacc_thread *thr = goacc_thread ();
struct gomp_device_descr *acc_dev = thr->dev;
+ if (acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
+ return;
+
gomp_mutex_lock (&acc_dev->lock);
n = lookup_host (acc_dev, h, s);
@@ -553,6 +588,9 @@ update_dev_host (int is_dev, void *h, size_t s)
struct goacc_thread *thr = goacc_thread ();
struct gomp_device_descr *acc_dev = thr->dev;
+ if (acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
+ return;
+
gomp_mutex_lock (&acc_dev->lock);
n = lookup_host (acc_dev, h, s);
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c
index f3b490a9508a3924474e7a7917a8d02577c4a705..d478ce2eef5799ad2cfec3c9274c463c916e92eb 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c
@@ -1,6 +1,4 @@
/* { dg-do run { target openacc_nvidia_accel_selected } } */
-/* <http://news.gmane.org/find-root.php?message_id=%3C87pp0aaksc.fsf%40kepler.schwinge.homeip.net%3E>.
- { dg-xfail-run-if "TODO" { *-*-* } } */
/* { dg-additional-options "-lcuda" } */
#include <openacc.h>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2-lib.c
new file mode 100644
index 0000000000000000000000000000000000000000..e1aa2c931ff047211bc79da61930f1117fbf2cd3
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2-lib.c
@@ -0,0 +1,185 @@
+/* This test is similar to data-2.c, but it uses acc_* library functions
+ to move data. */
+
+/* { dg-do run } */
+
+#include <stdlib.h>
+#include <assert.h>
+#include <openacc.h>
+
+int
+main (int argc, char **argv)
+{
+ int N = 128; //1024 * 1024;
+ float *a, *b, *c, *d, *e;
+ void *d_a, *d_b, *d_c, *d_d;
+ int i;
+ int nbytes;
+
+ nbytes = N * sizeof (float);
+
+ a = (float *) malloc (nbytes);
+ b = (float *) malloc (nbytes);
+ c = (float *) malloc (nbytes);
+ d = (float *) malloc (nbytes);
+ e = (float *) malloc (nbytes);
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 3.0;
+ b[i] = 0.0;
+ }
+
+ d_a = acc_copyin (a, nbytes);
+ d_b = acc_copyin (b, nbytes);
+ acc_copyin (&N, sizeof (int));
+
+#pragma acc parallel present (a[0:N], b[0:N], N) async wait
+#pragma acc loop
+ for (i = 0; i < N; i++)
+ b[i] = a[i];
+
+ acc_wait_all ();
+
+ acc_memcpy_from_device (a, d_a, nbytes);
+ acc_memcpy_from_device (b, d_b, nbytes);
+
+ for (i = 0; i < N; i++)
+ {
+ assert (a[i] == 3.0);
+ assert (b[i] == 3.0);
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 2.0;
+ b[i] = 0.0;
+ }
+
+ acc_update_device (a, nbytes);
+ acc_update_device (b, nbytes);
+
+#pragma acc parallel present (a[0:N], b[0:N], N) async (1)
+#pragma acc loop
+ for (i = 0; i < N; i++)
+ b[i] = a[i];
+
+ acc_memcpy_from_device (a, d_a, nbytes);
+ acc_memcpy_from_device (b, d_b, nbytes);
+
+ for (i = 0; i < N; i++)
+ {
+ assert (a[i] == 2.0);
+ assert (b[i] == 2.0);
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 3.0;
+ b[i] = 0.0;
+ c[i] = 0.0;
+ d[i] = 0.0;
+ }
+
+ acc_update_device (a, nbytes);
+ acc_update_device (b, nbytes);
+ d_c = acc_copyin (c, nbytes);
+ d_d = acc_copyin (d, nbytes);
+
+#pragma acc parallel present (a[0:N], b[0:N], N) async (1)
+#pragma acc loop
+ for (i = 0; i < N; i++)
+ b[i] = (a[i] * a[i] * a[i]) / a[i];
+
+#pragma acc parallel present (a[0:N], c[0:N], N) async (2)
+#pragma acc loop
+ for (i = 0; i < N; i++)
+ c[i] = (a[i] + a[i] + a[i] + a[i]) / a[i];
+
+#pragma acc parallel present (a[0:N], d[0:N], N) async (3)
+#pragma acc loop
+ for (i = 0; i < N; i++)
+ d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i];
+
+ acc_wait_all ();
+
+ acc_memcpy_from_device (a, d_a, nbytes);
+ acc_memcpy_from_device (b, d_b, nbytes);
+ acc_memcpy_from_device (c, d_c, nbytes);
+ acc_memcpy_from_device (d, d_d, nbytes);
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 3.0)
+ abort ();
+
+ if (b[i] != 9.0)
+ abort ();
+
+ if (c[i] != 4.0)
+ abort ();
+
+ if (d[i] != 1.0)
+ abort ();
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 2.0;
+ b[i] = 0.0;
+ c[i] = 0.0;
+ d[i] = 0.0;
+ e[i] = 0.0;
+ }
+
+ acc_update_device (a, nbytes);
+ acc_update_device (b, nbytes);
+ acc_update_device (c, nbytes);
+ acc_update_device (d, nbytes);
+ acc_copyin (e, nbytes);
+
+#pragma acc parallel present (a[0:N], b[0:N], N) async (1)
+ for (int ii = 0; ii < N; ii++)
+ b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
+
+#pragma acc parallel present (a[0:N], c[0:N], N) async (2)
+ for (int ii = 0; ii < N; ii++)
+ c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
+
+#pragma acc parallel present (a[0:N], d[0:N], N) async (3)
+ for (int ii = 0; ii < N; ii++)
+ d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
+
+#pragma acc parallel present (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N], N) \
+ async (4)
+ for (int ii = 0; ii < N; ii++)
+ e[ii] = a[ii] + b[ii] + c[ii] + d[ii];
+
+ acc_wait_all ();
+ acc_copyout (a, nbytes);
+ acc_copyout (b, nbytes);
+ acc_copyout (c, nbytes);
+ acc_copyout (d, nbytes);
+ acc_copyout (e, nbytes);
+ acc_delete (&N, sizeof (int));
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 2.0)
+ abort ();
+
+ if (b[i] != 4.0)
+ abort ();
+
+ if (c[i] != 4.0)
+ abort ();
+
+ if (d[i] != 1.0)
+ abort ();
+
+ if (e[i] != 11.0)
+ abort ();
+ }
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c
index f867a66e5b7f259135bd721c5860976941deeb40..c1c0825919d52da3bbe2f901d355579355ac6b8b 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c
@@ -1,3 +1,5 @@
+/* Test 'acc enter/exit data' regions. */
+
/* { dg-do run } */
#include <stdlib.h>
@@ -25,7 +27,7 @@ main (int argc, char **argv)
}
#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async
-#pragma acc parallel async wait
+#pragma acc parallel present (a[0:N], b[0:N]) async wait
#pragma acc loop
for (i = 0; i < N; i++)
b[i] = a[i];
@@ -49,7 +51,7 @@ main (int argc, char **argv)
}
#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async (1)
-#pragma acc parallel async (1)
+#pragma acc parallel present (a[0:N], b[0:N]) async (1)
#pragma acc loop
for (i = 0; i < N; i++)
b[i] = a[i];
@@ -76,17 +78,17 @@ main (int argc, char **argv)
#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (c[0:N]) copyin (d[0:N]) copyin (N) async (1)
-#pragma acc parallel async (1) wait (1)
+#pragma acc parallel present (a[0:N], b[0:N]) async (1) wait (1)
#pragma acc loop
for (i = 0; i < N; i++)
b[i] = (a[i] * a[i] * a[i]) / a[i];
-#pragma acc parallel async (2) wait (1)
+#pragma acc parallel present (a[0:N], c[0:N]) async (2) wait (1)
#pragma acc loop
for (i = 0; i < N; i++)
c[i] = (a[i] + a[i] + a[i] + a[i]) / a[i];
-#pragma acc parallel async (3) wait (1)
+#pragma acc parallel present (a[0:N], d[0:N]) async (3) wait (1)
#pragma acc loop
for (i = 0; i < N; i++)
d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i];
@@ -120,26 +122,27 @@ main (int argc, char **argv)
#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (c[0:N]) copyin (d[0:N]) copyin (e[0:N]) copyin (N) async (1)
-#pragma acc parallel async (1) wait (1)
+#pragma acc parallel present (a[0:N], b[0:N]) async (1) wait (1)
for (int ii = 0; ii < N; ii++)
b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
-#pragma acc parallel async (2) wait (1)
+#pragma acc parallel present (a[0:N], c[0:N]) async (2) wait (1)
for (int ii = 0; ii < N; ii++)
c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
-#pragma acc parallel async (3) wait (1)
+#pragma acc parallel present (a[0:N], d[0:N]) async (3) wait (1)
for (int ii = 0; ii < N; ii++)
d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
-#pragma acc parallel wait (1) async (4)
+#pragma acc parallel present (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) \
+ wait (1) async (4)
for (int ii = 0; ii < N; ii++)
e[ii] = a[ii] + b[ii] + c[ii] + d[ii];
-#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) copyout (d[0:N]) copyout (e[0:N]) wait (1, 2, 3, 4) async (1)
+#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) \
+ copyout (d[0:N]) copyout (e[0:N]) wait (1, 2, 3, 4) async (1)
#pragma acc wait (1)
-
for (i = 0; i < N; i++)
{
if (a[i] != 2.0)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c
index 747109f61363e0a81b939a8544ea636dbdc8a2f3..0bf706a1b5d42cf9983e61a28f926bc369554c6e 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c
@@ -1,3 +1,5 @@
+/* Test 'acc enter/exit data' regions with 'acc update'. */
+
/* { dg-do run } */
#include <stdlib.h>
@@ -25,7 +27,7 @@ main (int argc, char **argv)
}
#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async
-#pragma acc parallel async wait
+#pragma acc parallel present (a[0:N], b[0:N]) async wait
#pragma acc loop
for (i = 0; i < N; i++)
b[i] = a[i];
@@ -49,7 +51,7 @@ main (int argc, char **argv)
}
#pragma acc update device (a[0:N], b[0:N]) async (1)
-#pragma acc parallel async (1)
+#pragma acc parallel present (a[0:N], b[0:N]) async (1)
#pragma acc loop
for (i = 0; i < N; i++)
b[i] = a[i];
@@ -78,17 +80,17 @@ main (int argc, char **argv)
#pragma acc update device (b[0:N]) async (2)
#pragma acc enter data copyin (c[0:N], d[0:N]) async (3)
-#pragma acc parallel async (1) wait (1,2)
+#pragma acc parallel present (a[0:N], b[0:N]) async (1) wait (1,2)
#pragma acc loop
for (i = 0; i < N; i++)
b[i] = (a[i] * a[i] * a[i]) / a[i];
-#pragma acc parallel async (2) wait (1,3)
+#pragma acc parallel present (a[0:N], c[0:N]) async (2) wait (1,3)
#pragma acc loop
for (i = 0; i < N; i++)
c[i] = (a[i] + a[i] + a[i] + a[i]) / a[i];
-#pragma acc parallel async (3) wait (1,3)
+#pragma acc parallel present (a[0:N], d[0:N]) async (3) wait (1,3)
#pragma acc loop
for (i = 0; i < N; i++)
d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i];
@@ -123,27 +125,28 @@ main (int argc, char **argv)
#pragma acc update device (a[0:N], b[0:N], c[0:N], d[0:N]) async (1)
#pragma acc enter data copyin (e[0:N]) async (5)
-#pragma acc parallel async (1) wait (1)
+#pragma acc parallel present (a[0:N], b[0:N]) async (1) wait (1)
for (int ii = 0; ii < N; ii++)
b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
-#pragma acc parallel async (2) wait (1)
+#pragma acc parallel present (a[0:N], c[0:N]) async (2) wait (1)
for (int ii = 0; ii < N; ii++)
c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
-#pragma acc parallel async (3) wait (1)
+#pragma acc parallel present (a[0:N], d[0:N]) async (3) wait (1)
for (int ii = 0; ii < N; ii++)
d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
-#pragma acc parallel wait (1,5) async (4)
+#pragma acc parallel present (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) \
+ wait (1,5) async (4)
for (int ii = 0; ii < N; ii++)
e[ii] = a[ii] + b[ii] + c[ii] + d[ii];
-#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) copyout (d[0:N]) copyout (e[0:N]) wait (1, 2, 3, 4) async (1)
+#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) \
+ copyout (d[0:N]) copyout (e[0:N]) wait (1, 2, 3, 4) async (1)
#pragma acc exit data delete (N)
#pragma acc wait (1)
-
for (i = 0; i < N; i++)
{
if (a[i] != 2.0)
@@ -162,5 +165,11 @@ main (int argc, char **argv)
abort ();
}
+ free (a);
+ free (b);
+ free (c);
+ free (d);
+ free (e);
+
return 0;
}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/enter_exit-lib.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/enter_exit-lib.c
new file mode 100644
index 0000000000000000000000000000000000000000..b5b37b2893fbd2262733d2ca252c8108b8b38f09
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/enter_exit-lib.c
@@ -0,0 +1,70 @@
+/* Verify enter/exit data interoperablilty between pragmas and
+ acc library calls. */
+
+/* { dg-do run } */
+
+#include <stdlib.h>
+#include <assert.h>
+#include <openacc.h>
+
+int
+main ()
+{
+ int *p = (int *)malloc (sizeof (int));
+
+ /* Test 1: pragma input, library output. */
+
+#pragma acc enter data copyin (p[0:1])
+
+#pragma acc parallel present (p[0:1]) num_gangs (1)
+ {
+ p[0] = 1;
+ }
+
+ acc_copyout (p, sizeof (int));
+
+ assert (p[0] == 1);
+
+ /* Test 2: library input, pragma output. */
+
+ acc_copyin (p, sizeof (int));
+
+#pragma acc parallel present (p[0:1]) num_gangs (1)
+ {
+ p[0] = 2;
+ }
+
+#pragma acc exit data copyout (p[0:1])
+
+ assert (p[0] == 2);
+
+ /* Test 3: library input, library output. */
+
+ acc_copyin (p, sizeof (int));
+
+#pragma acc parallel present (p[0:1]) num_gangs (1)
+ {
+ p[0] = 3;
+ }
+
+ acc_copyout (p, sizeof (int));
+
+ assert (p[0] == 3);
+
+ /* Test 4: pragma input, pragma output. */
+
+#pragma acc enter data copyin (p[0:1])
+
+#pragma acc parallel present (p[0:1]) num_gangs (1)
+ {
+ p[0] = 3;
+ }
+
+#pragma acc exit data copyout (p[0:1])
+
+ assert (p[0] == 3);
+
+ free (p);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-13.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-13.c
index 7098ef3549c294b1fddf5fde039b817b4882e31e..d6655335e21571234cceee5761bcd312f81678cc 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-13.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-13.c
@@ -1,4 +1,6 @@
-/* { dg-do run } */
+/* Check acc_is_present and acc_delete. */
+
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
#include <stdlib.h>
#include <openacc.h>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-14.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-14.c
index a9632f786ee8389f94cab3a42657bb534426ddf5..ee21257c9a5d0da5542298360b82b68246202778 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-14.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-14.c
@@ -1,4 +1,6 @@
-/* { dg-do run } */
+/* Check acc_is_present. */
+
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
#include <stdlib.h>
#include <openacc.h>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-15.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-15.c
index 4f6a731be1dc41387aab94be2e59f53d7c0558f4..50c17011fe794b881a38ec15c89f830fa3e5365b 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-15.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-15.c
@@ -1,4 +1,6 @@
-/* { dg-do run } */
+/* Check acc_is_present and acc_copyout. */
+
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
#include <stdlib.h>
#include <openacc.h>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-16.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-16.c
index 28e4e5c65fae4a34c7a0e9539ccceb8759f5c98b..c81a78de26d869e27b0714b8acd03edc0212fa21 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-16.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-16.c
@@ -1,4 +1,6 @@
-/* { dg-do run } */
+/* Test if duplicate data mappings with acc_copy_in. */
+
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
#include <stdio.h>
#include <stdlib.h>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-17.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-17.c
index 7d1767e378734a8a31a3a9fbe39b84c0f0eb0a56..a3487e8f5bfc49ab513dde9942be4869b6a8152f 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-17.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-17.c
@@ -1,4 +1,7 @@
-/* { dg-do run } */
+/* Check acc_copyout failure with acc_device_nvidia. */
+
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
+
#include <stdio.h>
#include <stdlib.h>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-18.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-18.c
index 160b33c90d865bb592a23fc84f3128e360d30dfa..b686cc9481551d3ee449ca6a7df2c661904f9e9f 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-18.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-18.c
@@ -1,4 +1,6 @@
-/* { dg-do run } */
+/* Verify that acc_delete unregisters data mappings on the device. */
+
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
#include <stdio.h>
#include <stdlib.h>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-20.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-20.c
index 4f8e14c2218e0841981e0ea9a92db263dab29a3c..25ceb3a26af47febe93a8cfe29259d6e3f1ce5a9 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-20.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-20.c
@@ -1,4 +1,6 @@
-/* { dg-do run } */
+/* Exercise acc_copyin and acc_copyout on nvidia targets. */
+
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
#include <stdio.h>
#include <stdlib.h>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-21.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-21.c
index d9087008290dba6f46f14160432f88209e4656cc..b170f81229ceb17dd62ed83d79aba450b35957e2 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-21.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-21.c
@@ -1,4 +1,6 @@
-/* { dg-do run } */
+/* Exercise acc_copyin and acc_copyout on nvidia targets. */
+
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
#include <stdio.h>
#include <stdlib.h>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-22.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-22.c
index a6c0197b1abbfc2ad08a4504d51d6a7c6fbc590f..65ff440a528cfcc0c0656cfc0b0be0d3e50fabcc 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-22.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-22.c
@@ -1,4 +1,6 @@
-/* { dg-do run } */
+/* Exercise acc_copyin and acc_copyout on nvidia targets. */
+
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
#include <stdio.h>
#include <stdlib.h>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-23.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-23.c
index 2339dd6fa25ab54a965439cf558adac270d731ae..fd4dc5971a197c38bd91f2af79e970fdea142142 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-23.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-23.c
@@ -1,4 +1,6 @@
-/* { dg-do run } */
+/* Exercise acc_copyin and acc_copyout on nvidia targets. */
+
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
#include <stdio.h>
#include <stdlib.h>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-24.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-24.c
index d7de8e34f4cedfae0a336b4f25ad0f3013207fbf..09e2817f41dba372e5d0207fbbd9935fa8e3dea0 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-24.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-24.c
@@ -1,4 +1,6 @@
-/* { dg-do run } */
+/* Exercise acc_create, acc_is_present and acc_delete. */
+
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
#include <stdlib.h>
#include <openacc.h>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-25.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-25.c
index bb709d3d607571ceb2dbfb1f36ff27d47b888201..5f00ccb3885a3010da64bf713c036e342b6a41fd 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-25.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-25.c
@@ -1,4 +1,6 @@
-/* { dg-do run } */
+/* Exercise acc_create and acc_delete on nvidia targets. */
+
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
#include <stdio.h>
#include <stdlib.h>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-28.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-28.c
index 9304daa9f3f81ae1fc06ad057c4b253204fcba67..7a96ab26ebd15970d6897c68de3ff7e5a5344af2 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-28.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-28.c
@@ -1,4 +1,6 @@
-/* { dg-do run } */
+/* Exercise acc_delete with a NULL address on nvidia targets. */
+
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
#include <stdio.h>
#include <stdlib.h>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-29.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-29.c
index 92e38587f2d179857b8af701354b8c87894bab84..318a060f228d4df03994b01dbc0392cc39fc6251 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-29.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-29.c
@@ -1,4 +1,6 @@
-/* { dg-do run } */
+/* Exercise acc_delete with size zero on nvidia targets. */
+
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
#include <stdio.h>
#include <stdlib.h>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-30.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-30.c
index e81627da5c15ea5e249d209d7f19f20b0c68096e..9bc9ecc10689634df77ed63ec255abc83aaaaf26 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-30.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-30.c
@@ -1,4 +1,6 @@
-/* { dg-do run } */
+/* Exercise an invalid partial acc_delete on nvidia targets. */
+
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
#include <stdio.h>
#include <stdlib.h>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-34.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-34.c
index 031c7318e68626bd942dfe764a8a3c87527b9da7..a24916d13066c63a63ac249adb721d0b75de9372 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-34.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-34.c
@@ -1,4 +1,6 @@
-/* { dg-do run } */
+/* Exercise an invalid acc_present_or_create on nvidia targets. */
+
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
#include <stdio.h>
#include <stdlib.h>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-42.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-42.c
index de5d1c13cf89e8ecd1edea33c9c1c96ed5c62bb8..30b90d49c7b4ab1651d57ceea7ade63c964f7841 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-42.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-42.c
@@ -1,4 +1,6 @@
-/* { dg-do run } */
+/* Exercise acc_update_device on unmapped data on nvidia targets. */
+
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
#include <stdio.h>
#include <stdlib.h>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-43.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-43.c
index 0d593f056716ef93023240a3bcb5b07705e7fbcb..5db29124e9ec5bc48667aa14e4d3b62ae558129f 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-43.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-43.c
@@ -1,4 +1,6 @@
-/* { dg-do run } */
+/* Exercise acc_update_device with a NULL data address on nvidia targets. */
+
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
#include <stdio.h>
#include <stdlib.h>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-44.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-44.c
index e98ecc42794aaae386c1346ef3fe3e3ee781b83e..8bbf016a19192fda5ac40c81739b16ba4866bed5 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-44.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-44.c
@@ -1,4 +1,6 @@
-/* { dg-do run } */
+/* Exercise acc_update_device with size zero data on nvidia targets. */
+
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
#include <stdio.h>
#include <stdlib.h>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-47.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-47.c
index f26fc3318f606b1df50290a5e9de2b90d64b8de4..c2140429cb1edd115d792baa44397493c7a5b70f 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-47.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-47.c
@@ -1,4 +1,6 @@
-/* { dg-do run } */
+/* Exercise acc_update_self with a NULL data mapping on nvidia targets. */
+
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
#include <stdio.h>
#include <string.h>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-48.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-48.c
index 253ce59cefad37f3a92f5e3c7fdb3bdd590d7d58..afa137ff098bf49eeb66ca24dc809a9f8cf0ab23 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-48.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-48.c
@@ -1,4 +1,6 @@
-/* { dg-do run } */
+/* Exercise acc_update_self with a size zero data mapping on nvidia targets. */
+
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
#include <stdio.h>
#include <string.h>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-52.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-52.c
index cfbb07745327044f3e7aaac70cd9ca478dcb60a0..25c70c226f4be60531df207e9829b299958f5972 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-52.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-52.c
@@ -1,4 +1,6 @@
-/* { dg-do run } */
+/* Exercise acc_map_data with a NULL data mapping on nvidia targets. */
+
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
#include <stdio.h>
#include <stdlib.h>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-53.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-53.c
index 5de376d0328018adc69501990bb22371728fef5d..a8ee7df629c7052b27302e0f1efb2384886098d8 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-53.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-53.c
@@ -1,4 +1,6 @@
-/* { dg-do run } */
+/* Exercise acc_map_data with a NULL data mapping on nvidia targets. */
+
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
#include <stdio.h>
#include <stdlib.h>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-54.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-54.c
index 3e621c3c0588c6747b4f757e93f7b221591a5ae5..fc221f471167fdb7ac8d82d42bdc8a1b04d4c5d8 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-54.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-54.c
@@ -1,4 +1,6 @@
-/* { dg-do run } */
+/* Exercise acc_map_data with data size of zero on nvidia targets. */
+
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
#include <stdio.h>
#include <stdlib.h>