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

Fix oacc kernels default mapping for scalars

Browse Source 
2015-12-02  Tom de Vries  <tom@codesourcery.com>

	* gimplify.c (enum gimplify_omp_var_data): Add enum value
	GOVD_MAP_FORCE.
	(oacc_default_clause): Fix default for scalars in oacc kernels.
	(gimplify_adjust_omp_clauses_1): Handle GOVD_MAP_FORCE.

	* c-c++-common/goacc/kernels-default-2.c: New test.
	* c-c++-common/goacc/kernels-default.c: New test.

From-SVN: r231183
This commit is contained in:
Tom de Vries 2015-12-02 15:48:45 +00:00 committed by Tom de Vries
parent 86938de6f6
commit db0f1c7a90
5 changed files with 57 additions and 5 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

7
gcc/ChangeLog
View File

@ -1,3 +1,10 @@
2015-12-02 Tom de Vries <tom@codesourcery.com>
* gimplify.c (enum gimplify_omp_var_data): Add enum value
GOVD_MAP_FORCE.
(oacc_default_clause): Fix default for scalars in oacc kernels.
(gimplify_adjust_omp_clauses_1): Handle GOVD_MAP_FORCE.
2015-12-02 Tom de Vries <tom@codesourcery.com>
* omp-low.c (install_var_field, scan_sharing_clauses): Add and handle

19
gcc/gimplify.c
View File

@ -90,6 +90,9 @@ enum gimplify_omp_var_data
/* Flag for shared vars that are or might be stored to in the region. */
GOVD_WRITTEN = 131072,
/* Flag for GOVD_MAP, if it is a forced mapping. */
GOVD_MAP_FORCE = 262144,
GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
| GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
| GOVD_LOCAL)
@ -5980,8 +5983,12 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
gcc_unreachable ();
case ORT_ACC_KERNELS:
/* Everything under kernels are default 'present_or_copy'. */
/* Scalars are default 'copy' under kernels, non-scalars are default
'present_or_copy'. */
flags |= GOVD_MAP;
if (!AGGREGATE_TYPE_P (TREE_TYPE (decl)))
flags |= GOVD_MAP_FORCE;
rkind = "kernels";
break;
@ -7640,10 +7647,12 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
}
else if (code == OMP_CLAUSE_MAP)
{
OMP_CLAUSE_SET_MAP_KIND (clause,
flags & GOVD_MAP_TO_ONLY
? GOMP_MAP_TO
: GOMP_MAP_TOFROM);
int kind = (flags & GOVD_MAP_TO_ONLY
? GOMP_MAP_TO
: GOMP_MAP_TOFROM);
if (flags & GOVD_MAP_FORCE)
kind |= GOMP_MAP_FLAG_FORCE;
OMP_CLAUSE_SET_MAP_KIND (clause, kind);
if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
{

5
gcc/testsuite/ChangeLog
View File

@ -1,3 +1,8 @@
2015-12-02 Tom de Vries <tom@codesourcery.com>
* c-c++-common/goacc/kernels-default-2.c: New test.
* c-c++-common/goacc/kernels-default.c: New test.
2015-12-02 Tom de Vries <tom@codesourcery.com>
* c-c++-common/goacc/kernels-alias-2.c: New test.

17
gcc/testsuite/c-c++-common/goacc/kernels-default-2.c Normal file
View File

@ -0,0 +1,17 @@
/* { dg-additional-options "-O2" } */
/* { dg-additional-options "-fdump-tree-gimple" } */
#define N 2
void
foo (void)
{
unsigned int a[N];
#pragma acc kernels
{
a[0]++;
}
}
/* { dg-final { scan-tree-dump-times "map\\(tofrom" 1 "gimple" } } */

14
gcc/testsuite/c-c++-common/goacc/kernels-default.c Normal file
View File

@ -0,0 +1,14 @@
/* { dg-additional-options "-O2" } */
/* { dg-additional-options "-fdump-tree-gimple" } */
void
foo (void)
{
unsigned int i;
#pragma acc kernels
{
i++;
}
}
/* { dg-final { scan-tree-dump-times "map\\(force_tofrom" 1 "gimple" } } */