OpenE2K/gcc
2
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

[PR92036] Add 'libgomp.oacc-c-c++-common/data-firstprivate-1.c'

Browse Source 
libgomp/
	PR middle-end/92036
	* testsuite/libgomp.oacc-c-c++-common/data-firstprivate-1.c: New
	file.

From-SVN: r276757
This commit is contained in:
Thomas Schwinge 2019-10-09 13:31:14 +02:00 committed by Thomas Schwinge
parent e3423d768d
commit 6bbead0c5a
2 changed files with 171 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
libgomp/ChangeLog
View File

@ -1,3 +1,9 @@
2019-10-09 Thomas Schwinge <thomas@codesourcery.com>
PR middle-end/92036
* testsuite/libgomp.oacc-c-c++-common/data-firstprivate-1.c: New
file.
2019-10-09 Tobias Burnus <tobias@codesourcery.com>
PR testsuite/91884

165
libgomp/testsuite/libgomp.oacc-c-c++-common/data-firstprivate-1.c Normal file
View File

@ -0,0 +1,165 @@
/* Test behavior of 'firstprivate' lexically vs. dynamically nested inside a
'data' region. */
#include <stdlib.h>
#define VERIFY(x) \
do { \
if (!(x)) \
abort (); \
} while (0);
/* This is basically and extended version of 't2' from 'firstprivate-1.c'. */
int lexically_nested_val = 2;
static void
lexically_nested ()
{
#pragma acc data \
copy (lexically_nested_val)
{
VERIFY (lexically_nested_val == 2);
#pragma acc parallel \
present (lexically_nested_val)
{
VERIFY (lexically_nested_val == 2);
/* This updates the device copy, or shared variable. */
lexically_nested_val = 7;
}
#if ACC_MEM_SHARED
VERIFY (lexically_nested_val == 7);
#else
VERIFY (lexically_nested_val == 2);
#endif
/* This only updates the local/shared variable, but not the device
copy. */
lexically_nested_val = 5;
#pragma acc parallel \
firstprivate (lexically_nested_val)
{
#if 1 /* Current behavior. */
/* The 'firstprivate' copy is initialized from the device copy, or
shared variable. */
# if ACC_MEM_SHARED
VERIFY (lexically_nested_val == 5);
# else
VERIFY (lexically_nested_val == 7);
# endif
#else /* Expected behavior per PR92036. */
/* The 'firstprivate' copy is initialized from the local thread. */
VERIFY (lexically_nested_val == 5);
#endif
/* This updates the 'firstprivate' copy only, but not the shared
variable. */
lexically_nested_val = 9;
}
VERIFY (lexically_nested_val == 5);
}
/* If not shared, the device copy has now been copied back. */
#if ACC_MEM_SHARED
VERIFY (lexically_nested_val == 5);
#else
VERIFY (lexically_nested_val == 7);
#endif
}
int dynamically_nested_val = 2;
/* Same as above, but compute construct 1 broken out, so no longer lexically
nested inside 'data' region. */
static void
dynamically_nested_compute_1 ()
{
#pragma acc parallel \
present (dynamically_nested_val)
{
VERIFY (dynamically_nested_val == 2);
/* This updates the device copy, or shared variable. */
dynamically_nested_val = 7;
}
}
/* Same as above, but compute construct 2 broken out, so no longer lexically
nested inside 'data' region. */
static void
dynamically_nested_compute_2 ()
{
#pragma acc parallel \
firstprivate (dynamically_nested_val)
{
#if 1 /* Current behavior. */
/* The 'firstprivate' copy is initialized from the device copy, or shared
variable. */
# if ACC_MEM_SHARED
VERIFY (dynamically_nested_val == 5);
# else
VERIFY (dynamically_nested_val == 7);
# endif
#else /* Expected behavior per PR92036. */
/* The 'firstprivate' copy is initialized from the local thread. */
VERIFY (dynamically_nested_val == 5);
#endif
/* This updates the 'firstprivate' copy only, but not the shared
variable. */
dynamically_nested_val = 9;
}
}
static void
dynamically_nested ()
{
#pragma acc data \
copy (dynamically_nested_val)
{
VERIFY (dynamically_nested_val == 2);
dynamically_nested_compute_1 ();
#if ACC_MEM_SHARED
VERIFY (dynamically_nested_val == 7);
#else
VERIFY (dynamically_nested_val == 2);
#endif
/* This only updates the local/shared variable, but not the device
copy. */
dynamically_nested_val = 5;
dynamically_nested_compute_2 ();
VERIFY (dynamically_nested_val == 5);
}
/* If not shared, the device copy has now been copied back. */
#if ACC_MEM_SHARED
VERIFY (dynamically_nested_val == 5);
#else
VERIFY (dynamically_nested_val == 7);
#endif
}
int
main()
{
lexically_nested ();
dynamically_nested ();
return 0;
}