Commit fffeedeb by Nathan Sidwell Committed by Nathan Sidwell

gimplify.c (oacc_default_clause): New.

	gcc/
	* gimplify.c (oacc_default_clause): New.
	(omp_notice_variable): Call it.

	gcc/testsuite/
	* c-c++-common/goacc/data-default-1.c: New.

	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/default-1.c: New.

Co-Authored-By: Cesar Philippidis <cesar@codesourcery.com>

From-SVN: r230256
parent 8339a33e
2015-11-12 Nathan Sidwell <nathan@codesourcery.com>
gcc/
* gimplify.c (oacc_default_clause): New.
(omp_notice_variable): Call it.
gcc/testsuite/
* c-c++-common/goacc/data-default-1.c: New.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/default-1.c: New.
2015-11-12 Ilya Enkovich <enkovich.gnu@gmail.com> 2015-11-12 Ilya Enkovich <enkovich.gnu@gmail.com>
PR tree-optimization/68305 PR tree-optimization/68305
...@@ -311,7 +323,7 @@ ...@@ -311,7 +323,7 @@
2015-11-11 Nathan Sidwell <nathan@codesourcery.com> 2015-11-11 Nathan Sidwell <nathan@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com> Cesar Philippidis <cesar@codesourcery.com>
* gcc/gimplify.c (enum omp_region_type): Add ORT_ACC, * gimplify.c (enum omp_region_type): Add ORT_ACC,
ORT_ACC_DATA, ORT_ACC_PARALLEL, ORT_ACC_KERNELS. Adjust ORT_NONE. ORT_ACC_DATA, ORT_ACC_PARALLEL, ORT_ACC_KERNELS. Adjust ORT_NONE.
(gimple_add_tmp_var): Add ORT_ACC checks. (gimple_add_tmp_var): Add ORT_ACC checks.
(gimplify_var_or_parm_decl): Likewise. (gimplify_var_or_parm_decl): Likewise.
...@@ -327,7 +339,7 @@ ...@@ -327,7 +339,7 @@
(gimplify_oacc_cache): Specify ORT_ACC. (gimplify_oacc_cache): Specify ORT_ACC.
(gimplify_omp_workshare): Adjust OpenACC region types. (gimplify_omp_workshare): Adjust OpenACC region types.
(gimplify_omp_target_update): Likewise. (gimplify_omp_target_update): Likewise.
* gcc/omp-low.c (scan_sharing_clauses): Remove Openacc * omp-low.c (scan_sharing_clauses): Remove Openacc
firstprivate sorry. firstprivate sorry.
(lower-rec_input_clauses): Don't handle openacc firstprivate (lower-rec_input_clauses): Don't handle openacc firstprivate
references here. references here.
...@@ -5900,6 +5900,60 @@ omp_default_clause (struct gimplify_omp_ctx *ctx, tree decl, ...@@ -5900,6 +5900,60 @@ omp_default_clause (struct gimplify_omp_ctx *ctx, tree decl,
return flags; return flags;
} }
/* Determine outer default flags for DECL mentioned in an OACC region
but not declared in an enclosing clause. */
static unsigned
oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
{
const char *rkind;
switch (ctx->region_type)
{
default:
gcc_unreachable ();
case ORT_ACC_KERNELS:
/* Everything under kernels are default 'present_or_copy'. */
flags |= GOVD_MAP;
rkind = "kernels";
break;
case ORT_ACC_PARALLEL:
{
tree type = TREE_TYPE (decl);
if (TREE_CODE (type) == REFERENCE_TYPE
|| POINTER_TYPE_P (type))
type = TREE_TYPE (type);
if (AGGREGATE_TYPE_P (type))
/* Aggregates default to 'present_or_copy'. */
flags |= GOVD_MAP;
else
/* Scalars default to 'firstprivate'. */
flags |= GOVD_FIRSTPRIVATE;
rkind = "parallel";
}
break;
}
if (DECL_ARTIFICIAL (decl))
; /* We can get compiler-generated decls, and should not complain
about them. */
else if (ctx->default_kind == OMP_CLAUSE_DEFAULT_NONE)
{
error ("%qE not specified in enclosing OpenACC %s construct",
DECL_NAME (lang_hooks.decls.omp_report_decl (decl)), rkind);
error_at (ctx->location, "enclosing OpenACC %s construct", rkind);
}
else
gcc_checking_assert (ctx->default_kind == OMP_CLAUSE_DEFAULT_SHARED);
return flags;
}
/* Record the fact that DECL was used within the OMP context CTX. /* Record the fact that DECL was used within the OMP context CTX.
IN_CODE is true when real code uses DECL, and false when we should IN_CODE is true when real code uses DECL, and false when we should
merely emit default(none) errors. Return true if DECL is going to merely emit default(none) errors. Return true if DECL is going to
...@@ -6023,7 +6077,12 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code) ...@@ -6023,7 +6077,12 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
nflags |= GOVD_MAP | GOVD_EXPLICIT; nflags |= GOVD_MAP | GOVD_EXPLICIT;
} }
else if (nflags == flags) else if (nflags == flags)
nflags |= GOVD_MAP; {
if ((ctx->region_type & ORT_ACC) != 0)
nflags = oacc_default_clause (ctx, decl, flags);
else
nflags |= GOVD_MAP;
}
} }
found_outer: found_outer:
omp_add_variable (ctx, decl, nflags); omp_add_variable (ctx, decl, nflags);
......
2015-11-12 Nathan Sidwell <nathan@codesourcery.com>
* c-c++-common/goacc/data-default-1.c: New.
2015-11-12 David Edelsohn <dje.gcc@gmail.com> 2015-11-12 David Edelsohn <dje.gcc@gmail.com>
* gcc.target/powerpc/pr67789.c: Skip on AIX and Darwin. * gcc.target/powerpc/pr67789.c: Skip on AIX and Darwin.
...@@ -82,10 +86,6 @@ ...@@ -82,10 +86,6 @@
* gcc.dg/tree-ssa/pr68234.c: New testcase. * gcc.dg/tree-ssa/pr68234.c: New testcase.
2015-11-10 Nathan Sidwell <nathan@codesourcery.com>
* gcc.dg/goacc/nvptx-opt-1.c: New test.
2015-11-10 Ilya Enkovich <enkovich.gnu@gmail.com> 2015-11-10 Ilya Enkovich <enkovich.gnu@gmail.com>
* gcc.target/i386/mask-pack.c: New test. * gcc.target/i386/mask-pack.c: New test.
......
/* { dg-do compile } */
int main ()
{
int n = 2;
int ary[2];
#pragma acc parallel default (none) /* { dg-message "parallel construct" 2 } */
{
ary[0] /* { dg-error "not specified in enclosing" } */
= n; /* { dg-error "not specified in enclosing" } */
}
#pragma acc kernels default (none) /* { dg-message "kernels construct" 2 } */
{
ary[0] /* { dg-error "not specified in enclosing" } */
= n; /* { dg-error "not specified in enclosing" } */
}
#pragma acc data copy (ary, n)
{
#pragma acc parallel default (none)
{
ary[0]
= n;
}
#pragma acc kernels default (none)
{
ary[0]
= n;
}
}
return 0;
}
2015-11-12 Nathan Sidwell <nathan@codesourcery.com>
* testsuite/libgomp.oacc-c-c++-common/default-1.c: New.
2015-11-1 Nathan Sidwell <nathan@codesourcery.com> 2015-11-1 Nathan Sidwell <nathan@codesourcery.com>
* testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c: New. * testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c: New.
......
/* { dg-do run } */
#include <openacc.h>
int test_parallel ()
{
int ok = 1;
int val = 2;
int ary[32];
int ondev = 0;
for (int i = 0; i < 32; i++)
ary[i] = ~0;
/* val defaults to firstprivate, ary defaults to copy. */
#pragma acc parallel num_gangs (32) copy (ok) copy(ondev)
{
ondev = acc_on_device (acc_device_not_host);
#pragma acc loop gang(static:1)
for (unsigned i = 0; i < 32; i++)
{
if (val != 2)
ok = 0;
val += i;
ary[i] = val;
}
}
if (ondev)
{
if (!ok)
return 1;
if (val != 2)
return 1;
for (int i = 0; i < 32; i++)
if (ary[i] != 2 + i)
return 1;
}
return 0;
}
int test_kernels ()
{
int val = 2;
int ary[32];
int ondev = 0;
for (int i = 0; i < 32; i++)
ary[i] = ~0;
/* val defaults to copy, ary defaults to copy. */
#pragma acc kernels copy(ondev)
{
ondev = acc_on_device (acc_device_not_host);
#pragma acc loop
for (unsigned i = 0; i < 32; i++)
{
ary[i] = val;
val++;
}
}
if (ondev)
{
if (val != 2 + 32)
return 1;
for (int i = 0; i < 32; i++)
if (ary[i] != 2 + i)
return 1;
}
return 0;
}
int main ()
{
if (test_parallel ())
return 1;
if (test_kernels ())
return 1;
return 0;
}
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment