procyberian/gcc
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

c-parser.c (c_parser_oacc_declare): Add support for GOMP_MAP_FIRSTPRIVATE_POINTER.

Browse source 
gcc/c/
	* 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.

	gcc/cp/
	* 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.

	gcc/
	* 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.

	gcc/testsuite/
	* 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.

	libgomp/
	* 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.

From-SVN: r236678
This commit is contained in:
Cesar Philippidis 2016-05-24 15:54:21 -07:00 • committed by Cesar Philippidis
parent 4bfc9db7e6
commit e46c777050
55 changed files with 722 additions and 113 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

8
gcc/ChangeLog
View file

@ -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

12
gcc/c/ChangeLog
View file

@ -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.

1
gcc/c/c-parser.c
View file

@ -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:

59
gcc/c/c-typeck.c
View file

@ -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

13
gcc/cp/ChangeLog
View file

@ -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

1
gcc/cp/parser.c
View file

@ -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:

72
gcc/cp/semantics.c
View file

@ -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

18
gcc/gimplify.c
View file

@ -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)

18
gcc/testsuite/ChangeLog
View file

@ -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

8
gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c
View file

@ -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" } */
;
}

8
gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
View file

@ -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" } */
;
}

2
gcc/testsuite/c-c++-common/goacc/kernels-alias-3.c
View file

@ -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" } } */

2
gcc/testsuite/c-c++-common/goacc/kernels-alias-4.c
View file

@ -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" } } */

2
gcc/testsuite/c-c++-common/goacc/kernels-alias-5.c
View file

@ -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" } } */

4
gcc/testsuite/c-c++-common/goacc/kernels-alias-8.c
View file

@ -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" } } */

5
gcc/testsuite/c-c++-common/goacc/kernels-alias-ipa-pta-3.c
View file

@ -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 } } } */

2
gcc/testsuite/c-c++-common/goacc/pcopy.c
View file

@ -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" } } */

2
gcc/testsuite/c-c++-common/goacc/pcopyin.c
View file

@ -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" } } */

2
gcc/testsuite/c-c++-common/goacc/pcopyout.c
View file

@ -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" } } */

2
gcc/testsuite/c-c++-common/goacc/pcreate.c
View file

@ -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" } } */

48
gcc/testsuite/c-c++-common/goacc/pr70688.c Normal file
View file

@ -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;
}

2
gcc/testsuite/c-c++-common/goacc/present-1.c
View file

@ -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" } } */

4
gcc/testsuite/c-c++-common/goacc/reduction-5.c
View file

@ -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;

39
gcc/testsuite/g++.dg/goacc/data-1.C Normal file
View file

@ -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 } */

45
libgomp/ChangeLog
View file

@ -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.

46
libgomp/oacc-mem.c
View file

@ -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);

2
libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c
View file

@ -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>

185
libgomp/testsuite/libgomp.oacc-c-c++-common/data-2-lib.c Normal file
View file

@ -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;
}

25
libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c
View file

@ -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)

31
libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c
View file

@ -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;
}

70
libgomp/testsuite/libgomp.oacc-c-c++-common/enter_exit-lib.c Normal file
View file

@ -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;
}

4
libgomp/testsuite/libgomp.oacc-c-c++-common/lib-13.c
View file

@ -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>

4
libgomp/testsuite/libgomp.oacc-c-c++-common/lib-14.c
View file

@ -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>

4
libgomp/testsuite/libgomp.oacc-c-c++-common/lib-15.c
View file

@ -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>

4
libgomp/testsuite/libgomp.oacc-c-c++-common/lib-16.c
View file

@ -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>

5
libgomp/testsuite/libgomp.oacc-c-c++-common/lib-17.c
View file

@ -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>

4
libgomp/testsuite/libgomp.oacc-c-c++-common/lib-18.c
View file

@ -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>

4
libgomp/testsuite/libgomp.oacc-c-c++-common/lib-20.c
View file

@ -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>

4
libgomp/testsuite/libgomp.oacc-c-c++-common/lib-21.c
View file

@ -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>

4
libgomp/testsuite/libgomp.oacc-c-c++-common/lib-22.c
View file

@ -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>

4
libgomp/testsuite/libgomp.oacc-c-c++-common/lib-23.c
View file

@ -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>

4
libgomp/testsuite/libgomp.oacc-c-c++-common/lib-24.c
View file

@ -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>

4
libgomp/testsuite/libgomp.oacc-c-c++-common/lib-25.c
View file

@ -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>

4
libgomp/testsuite/libgomp.oacc-c-c++-common/lib-28.c
View file

@ -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>

4
libgomp/testsuite/libgomp.oacc-c-c++-common/lib-29.c
View file

@ -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>

4
libgomp/testsuite/libgomp.oacc-c-c++-common/lib-30.c
View file

@ -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>

4
libgomp/testsuite/libgomp.oacc-c-c++-common/lib-34.c
View file

@ -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>

4
libgomp/testsuite/libgomp.oacc-c-c++-common/lib-42.c
View file

@ -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>

4
libgomp/testsuite/libgomp.oacc-c-c++-common/lib-43.c
View file

@ -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>

4
libgomp/testsuite/libgomp.oacc-c-c++-common/lib-44.c
View file

@ -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>

4
libgomp/testsuite/libgomp.oacc-c-c++-common/lib-47.c
View file

@ -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>

4
libgomp/testsuite/libgomp.oacc-c-c++-common/lib-48.c
View file

@ -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>

4
libgomp/testsuite/libgomp.oacc-c-c++-common/lib-52.c
View file

@ -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>

4
libgomp/testsuite/libgomp.oacc-c-c++-common/lib-53.c
View file

@ -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>

4
libgomp/testsuite/libgomp.oacc-c-c++-common/lib-54.c
View file

@ -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>