git
/
gcc
mirror of git://gcc.gnu.org/git/gcc.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

gimplify.c (enum omp_region_type): Add ORT_ACC, ORT_ACC_DATA, ORT_ACC_PARALLEL, ORT_ACC_KERNELS. 

gcc/
	* gcc/gimplify.c (enum  omp_region_type): Add ORT_ACC,
	ORT_ACC_DATA, ORT_ACC_PARALLEL, ORT_ACC_KERNELS.  Adjust ORT_NONE.
	(gimple_add_tmp_var): Add ORT_ACC checks.
	(gimplify_var_or_parm_decl): Likewise.
	(omp_firstprivatize_variable): Likewise. Use ORT_TARGET_DATA as a
	mask.
	(omp_add_variable): Look in outer contexts for openacc and allow
	reductions with other sharing. Add ORT_ACC and ORT_TARGET_DATA
	checks.
	(omp_notice_variable, omp_is_private, omp_check_private): Add
	ORT_ACC checks.
	(gimplify_scan_omp_clauses: Treat ORT_ACC as ORT_WORKSHARE.
	Permit private openacc reductions.
	(gimplify_oacc_cache): Specify ORT_ACC.
	(gimplify_omp_workshare): Adjust OpenACC region types.
	(gimplify_omp_target_update): Likewise.
	* gcc/omp-low.c (scan_sharing_clauses): Remove Openacc
	firstprivate sorry.
	(lower-rec_input_clauses): Don't handle openacc firstprivate
	references here.
	(lower_omp_target): Emit initializers for openacc firstprivate vars.

	gcc/testsuite/
	* gfortran.dg/goacc/private-3.f95: Remove xfail.
	* gfortran.dg/goacc/combined_loop.f90: Remove xfail.

	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c: Remove xfail.
	* testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c: Remove xfail.
	* testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c: New.
	* testsuite/libgomp.oacc-c-c++-common/firstprivate-2.c: New.

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

From-SVN: r230169
This commit is contained in:
Nathan Sidwell 2015-11-11 14:24:09 +00:00 committed by Nathan Sidwell
parent 7700cd858f
commit 182190f2b3
11 changed files with 278 additions and 72 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

26
gcc/ChangeLog
View File

@ -1,3 +1,29 @@
2015-11-11 Nathan Sidwell <nathan@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
gcc/
* gcc/gimplify.c (enum omp_region_type): Add ORT_ACC,
ORT_ACC_DATA, ORT_ACC_PARALLEL, ORT_ACC_KERNELS. Adjust ORT_NONE.
(gimple_add_tmp_var): Add ORT_ACC checks.
(gimplify_var_or_parm_decl): Likewise.
(omp_firstprivatize_variable): Likewise. Use ORT_TARGET_DATA as a
mask.
(omp_add_variable): Look in outer contexts for openacc and allow
reductions with other sharing. Add ORT_ACC and ORT_TARGET_DATA
checks.
(omp_notice_variable, omp_is_private, omp_check_private): Add
ORT_ACC checks.
(gimplify_scan_omp_clauses: Treat ORT_ACC as ORT_WORKSHARE.
Permit private openacc reductions.
(gimplify_oacc_cache): Specify ORT_ACC.
(gimplify_omp_workshare): Adjust OpenACC region types.
(gimplify_omp_target_update): Likewise.
* gcc/omp-low.c (scan_sharing_clauses): Remove Openacc
firstprivate sorry.
(lower-rec_input_clauses): Don't handle openacc firstprivate
references here.
(lower_omp_target): Emit initializers for openacc firstprivate vars.
2015-11-11 Eric Botcazou <ebotcazou@adacore.com> 2015-11-11 Eric Botcazou <ebotcazou@adacore.com>
PR target/67265 PR target/67265

141
gcc/gimplify.c
View File

@ -95,22 +95,34 @@ enum gimplify_omp_var_data
enum omp_region_type enum omp_region_type
{ {
ORT_WORKSHARE = 0, ORT_WORKSHARE = 0x00,
ORT_SIMD = 1, ORT_SIMD = 0x01,
ORT_PARALLEL = 2,
ORT_COMBINED_PARALLEL = 3, ORT_PARALLEL = 0x02,
ORT_TASK = 4, ORT_COMBINED_PARALLEL = 0x03,
ORT_UNTIED_TASK = 5,
ORT_TEAMS = 8, ORT_TASK = 0x04,
ORT_COMBINED_TEAMS = 9, ORT_UNTIED_TASK = 0x05,
ORT_TEAMS = 0x08,
ORT_COMBINED_TEAMS = 0x09,
/* Data region. */ /* Data region. */
ORT_TARGET_DATA = 16, ORT_TARGET_DATA = 0x10,
/* Data region with offloading. */ /* Data region with offloading. */
ORT_TARGET = 32, ORT_TARGET = 0x20,
ORT_COMBINED_TARGET = 33, ORT_COMBINED_TARGET = 0x21,
/* OpenACC variants. */
ORT_ACC = 0x40, /* A generic OpenACC region. */
ORT_ACC_DATA = ORT_ACC | ORT_TARGET_DATA, /* Data construct. */
ORT_ACC_PARALLEL = ORT_ACC | ORT_TARGET, /* Parallel construct */
ORT_ACC_KERNELS = ORT_ACC | ORT_TARGET | 0x80, /* Kernels construct. */
/* Dummy OpenMP region, used to disable expansion of /* Dummy OpenMP region, used to disable expansion of
DECL_VALUE_EXPRs in taskloop pre body. */ DECL_VALUE_EXPRs in taskloop pre body. */
ORT_NONE = 64 ORT_NONE = 0x100
}; };
/* Gimplify hashtable helper. */ /* Gimplify hashtable helper. */
@ -689,7 +701,8 @@ gimple_add_tmp_var (tree tmp)
struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp; struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
while (ctx while (ctx
&& (ctx->region_type == ORT_WORKSHARE && (ctx->region_type == ORT_WORKSHARE
|| ctx->region_type == ORT_SIMD)) || ctx->region_type == ORT_SIMD
|| ctx->region_type == ORT_ACC))
ctx = ctx->outer_context; ctx = ctx->outer_context;
if (ctx) if (ctx)
omp_add_variable (ctx, tmp, GOVD_LOCAL | GOVD_SEEN); omp_add_variable (ctx, tmp, GOVD_LOCAL | GOVD_SEEN);
@ -1804,7 +1817,8 @@ gimplify_var_or_parm_decl (tree *expr_p)
struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp; struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
while (ctx while (ctx
&& (ctx->region_type == ORT_WORKSHARE && (ctx->region_type == ORT_WORKSHARE
|| ctx->region_type == ORT_SIMD)) || ctx->region_type == ORT_SIMD
|| ctx->region_type == ORT_ACC))
ctx = ctx->outer_context; ctx = ctx->outer_context;
if (!ctx && !nonlocal_vlas->add (decl)) if (!ctx && !nonlocal_vlas->add (decl))
{ {
@ -5579,7 +5593,8 @@ omp_firstprivatize_variable (struct gimplify_omp_ctx *ctx, tree decl)
} }
else if (ctx->region_type != ORT_WORKSHARE else if (ctx->region_type != ORT_WORKSHARE
&& ctx->region_type != ORT_SIMD && ctx->region_type != ORT_SIMD
&& ctx->region_type != ORT_TARGET_DATA) && ctx->region_type != ORT_ACC
&& !(ctx->region_type & ORT_TARGET_DATA))
omp_add_variable (ctx, decl, GOVD_FIRSTPRIVATE); omp_add_variable (ctx, decl, GOVD_FIRSTPRIVATE);
ctx = ctx->outer_context; ctx = ctx->outer_context;
@ -5667,11 +5682,13 @@ omp_add_variable (struct gimplify_omp_ctx *ctx, tree decl, unsigned int flags)
/* We shouldn't be re-adding the decl with the same data /* We shouldn't be re-adding the decl with the same data
sharing class. */ sharing class. */
gcc_assert ((n->value & GOVD_DATA_SHARE_CLASS & flags) == 0); gcc_assert ((n->value & GOVD_DATA_SHARE_CLASS & flags) == 0);
/* The only combination of data sharing classes we should see is
FIRSTPRIVATE and LASTPRIVATE. */
nflags = n->value | flags; nflags = n->value | flags;
gcc_assert ((nflags & GOVD_DATA_SHARE_CLASS) /* The only combination of data sharing classes we should see is
== (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE) FIRSTPRIVATE and LASTPRIVATE. However, OpenACC permits
reduction variables to be used in data sharing clauses. */
gcc_assert ((ctx->region_type & ORT_ACC) != 0
|| ((nflags & GOVD_DATA_SHARE_CLASS)
== (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE))
|| (flags & GOVD_DATA_SHARE_CLASS) == 0); || (flags & GOVD_DATA_SHARE_CLASS) == 0);
n->value = nflags; n->value = nflags;
return; return;
@ -5968,20 +5985,47 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
else if (is_scalar) else if (is_scalar)
nflags |= GOVD_FIRSTPRIVATE; nflags |= GOVD_FIRSTPRIVATE;
} }
tree type = TREE_TYPE (decl);
if (nflags == flags struct gimplify_omp_ctx *octx = ctx->outer_context;
&& gimplify_omp_ctxp->target_firstprivatize_array_bases if ((ctx->region_type & ORT_ACC) && octx)
&& lang_hooks.decls.omp_privatize_by_reference (decl))
type = TREE_TYPE (type);
if (nflags == flags
&& !lang_hooks.types.omp_mappable_type (type))
{ {
error ("%qD referenced in target region does not have " /* Look in outer OpenACC contexts, to see if there's a
"a mappable type", decl); data attribute for this variable. */
nflags |= GOVD_MAP | GOVD_EXPLICIT; omp_notice_variable (octx, decl, in_code);
for (; octx; octx = octx->outer_context)
{
if (!(octx->region_type & (ORT_TARGET_DATA | ORT_TARGET)))
break;
splay_tree_node n2
= splay_tree_lookup (octx->variables,
(splay_tree_key) decl);
if (n2)
{
nflags |= GOVD_MAP;
goto found_outer;
}
}
} }
else if (nflags == flags)
nflags |= GOVD_MAP; {
tree type = TREE_TYPE (decl);
if (nflags == flags
&& gimplify_omp_ctxp->target_firstprivatize_array_bases
&& lang_hooks.decls.omp_privatize_by_reference (decl))
type = TREE_TYPE (type);
if (nflags == flags
&& !lang_hooks.types.omp_mappable_type (type))
{
error ("%qD referenced in target region does not have "
"a mappable type", decl);
nflags |= GOVD_MAP | GOVD_EXPLICIT;
}
else if (nflags == flags)
nflags |= GOVD_MAP;
}
found_outer:
omp_add_variable (ctx, decl, nflags); omp_add_variable (ctx, decl, nflags);
} }
else else
@ -5998,7 +6042,8 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
{ {
if (ctx->region_type == ORT_WORKSHARE if (ctx->region_type == ORT_WORKSHARE
|| ctx->region_type == ORT_SIMD || ctx->region_type == ORT_SIMD
|| ctx->region_type == ORT_TARGET_DATA) || ctx->region_type == ORT_ACC
|| (ctx->region_type & ORT_TARGET_DATA) != 0)
goto do_outer; goto do_outer;
flags = omp_default_clause (ctx, decl, in_code, flags); flags = omp_default_clause (ctx, decl, in_code, flags);
@ -6112,7 +6157,8 @@ omp_is_private (struct gimplify_omp_ctx *ctx, tree decl, int simd)
} }
if (ctx->region_type != ORT_WORKSHARE if (ctx->region_type != ORT_WORKSHARE
&& ctx->region_type != ORT_SIMD) && ctx->region_type != ORT_SIMD
&& ctx->region_type != ORT_ACC)
return false; return false;
else if (ctx->outer_context) else if (ctx->outer_context)
return omp_is_private (ctx->outer_context, decl, simd); return omp_is_private (ctx->outer_context, decl, simd);
@ -6168,7 +6214,8 @@ omp_check_private (struct gimplify_omp_ctx *ctx, tree decl, bool copyprivate)
} }
} }
while (ctx->region_type == ORT_WORKSHARE while (ctx->region_type == ORT_WORKSHARE
|| ctx->region_type == ORT_SIMD); || ctx->region_type == ORT_SIMD
|| ctx->region_type == ORT_ACC);
return false; return false;
} }
@ -6311,7 +6358,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
omp_notice_variable (outer_ctx->outer_context, decl, true); omp_notice_variable (outer_ctx->outer_context, decl, true);
} }
else if (outer_ctx else if (outer_ctx
&& outer_ctx->region_type == ORT_WORKSHARE && (outer_ctx->region_type == ORT_WORKSHARE
|| outer_ctx->region_type == ORT_ACC)
&& outer_ctx->combined_loop && outer_ctx->combined_loop
&& splay_tree_lookup (outer_ctx->variables, && splay_tree_lookup (outer_ctx->variables,
(splay_tree_key) decl) == NULL (splay_tree_key) decl) == NULL
@ -6335,7 +6383,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
goto do_add; goto do_add;
case OMP_CLAUSE_REDUCTION: case OMP_CLAUSE_REDUCTION:
flags = GOVD_REDUCTION | GOVD_SEEN | GOVD_EXPLICIT; flags = GOVD_REDUCTION | GOVD_SEEN | GOVD_EXPLICIT;
check_non_private = "reduction"; /* OpenACC permits reductions on private variables. */
if (!(region_type & ORT_ACC))
check_non_private = "reduction";
decl = OMP_CLAUSE_DECL (c); decl = OMP_CLAUSE_DECL (c);
if (TREE_CODE (decl) == MEM_REF) if (TREE_CODE (decl) == MEM_REF)
{ {
@ -7704,7 +7754,7 @@ gimplify_oacc_cache (tree *expr_p, gimple_seq *pre_p)
{ {
tree expr = *expr_p; tree expr = *expr_p;
gimplify_scan_omp_clauses (&OACC_CACHE_CLAUSES (expr), pre_p, ORT_WORKSHARE, gimplify_scan_omp_clauses (&OACC_CACHE_CLAUSES (expr), pre_p, ORT_ACC,
OACC_CACHE); OACC_CACHE);
gimplify_adjust_omp_clauses (pre_p, &OACC_CACHE_CLAUSES (expr), OACC_CACHE); gimplify_adjust_omp_clauses (pre_p, &OACC_CACHE_CLAUSES (expr), OACC_CACHE);
@ -7833,7 +7883,9 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
case OMP_FOR: case OMP_FOR:
case CILK_FOR: case CILK_FOR:
case OMP_DISTRIBUTE: case OMP_DISTRIBUTE:
break;
case OACC_LOOP: case OACC_LOOP:
ort = ORT_ACC;
break; break;
case OMP_TASKLOOP: case OMP_TASKLOOP:
if (find_omp_clause (OMP_FOR_CLAUSES (for_stmt), OMP_CLAUSE_UNTIED)) if (find_omp_clause (OMP_FOR_CLAUSES (for_stmt), OMP_CLAUSE_UNTIED))
@ -8895,10 +8947,14 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
ort = OMP_TARGET_COMBINED (expr) ? ORT_COMBINED_TARGET : ORT_TARGET; ort = OMP_TARGET_COMBINED (expr) ? ORT_COMBINED_TARGET : ORT_TARGET;
break; break;
case OACC_KERNELS: case OACC_KERNELS:
ort = ORT_ACC_KERNELS;
break;
case OACC_PARALLEL: case OACC_PARALLEL:
ort = ORT_TARGET; ort = ORT_ACC_PARALLEL;
break; break;
case OACC_DATA: case OACC_DATA:
ort = ORT_ACC_DATA;
break;
case OMP_TARGET_DATA: case OMP_TARGET_DATA:
ort = ORT_TARGET_DATA; ort = ORT_TARGET_DATA;
break; break;
@ -8920,7 +8976,7 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
pop_gimplify_context (g); pop_gimplify_context (g);
else else
pop_gimplify_context (NULL); pop_gimplify_context (NULL);
if (ort == ORT_TARGET_DATA) if ((ort & ORT_TARGET_DATA) != 0)
{ {
enum built_in_function end_ix; enum built_in_function end_ix;
switch (TREE_CODE (expr)) switch (TREE_CODE (expr))
@ -8995,17 +9051,18 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
tree expr = *expr_p; tree expr = *expr_p;
int kind; int kind;
gomp_target *stmt; gomp_target *stmt;
enum omp_region_type ort = ORT_WORKSHARE;
switch (TREE_CODE (expr)) switch (TREE_CODE (expr))
{ {
case OACC_ENTER_DATA: case OACC_ENTER_DATA:
kind = GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA;
break;
case OACC_EXIT_DATA: case OACC_EXIT_DATA:
kind = GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA; kind = GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA;
ort = ORT_ACC;
break; break;
case OACC_UPDATE: case OACC_UPDATE:
kind = GF_OMP_TARGET_KIND_OACC_UPDATE; kind = GF_OMP_TARGET_KIND_OACC_UPDATE;
ort = ORT_ACC;
break; break;
case OMP_TARGET_UPDATE: case OMP_TARGET_UPDATE:
kind = GF_OMP_TARGET_KIND_UPDATE; kind = GF_OMP_TARGET_KIND_UPDATE;
@ -9020,7 +9077,7 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
gcc_unreachable (); gcc_unreachable ();
} }
gimplify_scan_omp_clauses (&OMP_STANDALONE_CLAUSES (expr), pre_p, gimplify_scan_omp_clauses (&OMP_STANDALONE_CLAUSES (expr), pre_p,
ORT_WORKSHARE, TREE_CODE (expr)); ort, TREE_CODE (expr));
gimplify_adjust_omp_clauses (pre_p, &OMP_STANDALONE_CLAUSES (expr), gimplify_adjust_omp_clauses (pre_p, &OMP_STANDALONE_CLAUSES (expr),
TREE_CODE (expr)); TREE_CODE (expr));
stmt = gimple_build_omp_target (NULL, kind, OMP_STANDALONE_CLAUSES (expr)); stmt = gimple_build_omp_target (NULL, kind, OMP_STANDALONE_CLAUSES (expr));

93
gcc/omp-low.c
View File

@ -1896,12 +1896,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
/* FALLTHRU */ /* FALLTHRU */
case OMP_CLAUSE_FIRSTPRIVATE: case OMP_CLAUSE_FIRSTPRIVATE:
if (is_gimple_omp_oacc (ctx->stmt))
{
sorry ("clause not supported yet");
break;
}
/* FALLTHRU */
case OMP_CLAUSE_LINEAR: case OMP_CLAUSE_LINEAR:
decl = OMP_CLAUSE_DECL (c); decl = OMP_CLAUSE_DECL (c);
do_private: do_private:
@ -2167,12 +2161,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
/* FALLTHRU */ /* FALLTHRU */
case OMP_CLAUSE_FIRSTPRIVATE: case OMP_CLAUSE_FIRSTPRIVATE:
if (is_gimple_omp_oacc (ctx->stmt))
{
sorry ("clause not supported yet");
break;
}
/* FALLTHRU */
case OMP_CLAUSE_PRIVATE: case OMP_CLAUSE_PRIVATE:
case OMP_CLAUSE_LINEAR: case OMP_CLAUSE_LINEAR:
case OMP_CLAUSE_IS_DEVICE_PTR: case OMP_CLAUSE_IS_DEVICE_PTR:
@ -4684,7 +4672,7 @@ lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist,
gimplify_assign (ptr, x, ilist); gimplify_assign (ptr, x, ilist);
} }
} }
else if (is_reference (var)) else if (is_reference (var) && !is_oacc_parallel (ctx))
{ {
/* For references that are being privatized for Fortran, /* For references that are being privatized for Fortran,
allocate new backing storage for the new pointer allocate new backing storage for the new pointer
@ -14911,7 +14899,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
tree child_fn, t, c; tree child_fn, t, c;
gomp_target *stmt = as_a <gomp_target *> (gsi_stmt (*gsi_p)); gomp_target *stmt = as_a <gomp_target *> (gsi_stmt (*gsi_p));
gbind *tgt_bind, *bind, *dep_bind = NULL; gbind *tgt_bind, *bind, *dep_bind = NULL;
gimple_seq tgt_body, olist, ilist, new_body; gimple_seq tgt_body, olist, ilist, fplist, new_body;
location_t loc = gimple_location (stmt); location_t loc = gimple_location (stmt);
bool offloaded, data_region; bool offloaded, data_region;
unsigned int map_cnt = 0; unsigned int map_cnt = 0;
@ -14963,6 +14951,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
child_fn = ctx->cb.dst_fn; child_fn = ctx->cb.dst_fn;
push_gimplify_context (); push_gimplify_context ();
fplist = NULL;
for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c)) for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
switch (OMP_CLAUSE_CODE (c)) switch (OMP_CLAUSE_CODE (c))
@ -15007,6 +14996,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
/* FALLTHRU */ /* FALLTHRU */
case OMP_CLAUSE_TO: case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM: case OMP_CLAUSE_FROM:
oacc_firstprivate:
var = OMP_CLAUSE_DECL (c); var = OMP_CLAUSE_DECL (c);
if (!DECL_P (var)) if (!DECL_P (var))
{ {
@ -15029,6 +15019,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
} }
if (offloaded if (offloaded
&& OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_REFERENCE)) || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
{ {
@ -15057,17 +15048,40 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
x = build_receiver_ref (var, true, ctx); x = build_receiver_ref (var, true, ctx);
tree new_var = lookup_decl (var, ctx); tree new_var = lookup_decl (var, ctx);
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c) && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
&& TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE) && TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
x = build_simple_mem_ref (x); x = build_simple_mem_ref (x);
SET_DECL_VALUE_EXPR (new_var, x); if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
DECL_HAS_VALUE_EXPR_P (new_var) = 1; {
gcc_assert (is_gimple_omp_oacc (ctx->stmt));
if (is_reference (new_var))
{
/* Create a local object to hold the instance
value. */
tree type = TREE_TYPE (TREE_TYPE (new_var));
const char *id = IDENTIFIER_POINTER (DECL_NAME (new_var));
tree inst = create_tmp_var (type, id);
gimplify_assign (inst, fold_indirect_ref (x), &fplist);
x = build_fold_addr_expr (inst);
}
gimplify_assign (new_var, x, &fplist);
}
else if (DECL_P (new_var))
{
SET_DECL_VALUE_EXPR (new_var, x);
DECL_HAS_VALUE_EXPR_P (new_var) = 1;
}
else
gcc_unreachable ();
} }
map_cnt++; map_cnt++;
break; break;
case OMP_CLAUSE_FIRSTPRIVATE: case OMP_CLAUSE_FIRSTPRIVATE:
if (is_oacc_parallel (ctx))
goto oacc_firstprivate;
map_cnt++; map_cnt++;
var = OMP_CLAUSE_DECL (c); var = OMP_CLAUSE_DECL (c);
if (!is_reference (var) if (!is_reference (var)
@ -15092,6 +15106,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
break; break;
case OMP_CLAUSE_PRIVATE: case OMP_CLAUSE_PRIVATE:
if (is_gimple_omp_oacc (ctx->stmt))
break;
var = OMP_CLAUSE_DECL (c); var = OMP_CLAUSE_DECL (c);
if (is_variable_sized (var)) if (is_variable_sized (var))
{ {
@ -15195,9 +15211,11 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
default: default:
break; break;
case OMP_CLAUSE_MAP: case OMP_CLAUSE_MAP:
case OMP_CLAUSE_TO: case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM: case OMP_CLAUSE_FROM:
oacc_firstprivate_map:
nc = c; nc = c;
ovar = OMP_CLAUSE_DECL (c); ovar = OMP_CLAUSE_DECL (c);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
@ -15248,9 +15266,9 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
x = build_sender_ref (ovar, ctx); x = build_sender_ref (ovar, ctx);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c) && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
&& TREE_CODE (TREE_TYPE (ovar)) == ARRAY_TYPE) && TREE_CODE (TREE_TYPE (ovar)) == ARRAY_TYPE)
{ {
gcc_assert (offloaded); gcc_assert (offloaded);
tree avar tree avar
@ -15261,6 +15279,15 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
avar = build_fold_addr_expr (avar); avar = build_fold_addr_expr (avar);
gimplify_assign (x, avar, &ilist); gimplify_assign (x, avar, &ilist);
} }
else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
{
gcc_assert (is_gimple_omp_oacc (ctx->stmt));
if (!is_reference (var))
var = build_fold_addr_expr (var);
else
talign = TYPE_ALIGN_UNIT (TREE_TYPE (TREE_TYPE (ovar)));
gimplify_assign (x, var, &ilist);
}
else if (is_gimple_reg (var)) else if (is_gimple_reg (var))
{ {
gcc_assert (offloaded); gcc_assert (offloaded);
@ -15289,7 +15316,17 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
gimplify_assign (x, var, &ilist); gimplify_assign (x, var, &ilist);
} }
} }
s = OMP_CLAUSE_SIZE (c); s = NULL_TREE;
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
{
gcc_checking_assert (is_gimple_omp_oacc (ctx->stmt));
s = TREE_TYPE (ovar);
if (TREE_CODE (s) == REFERENCE_TYPE)
s = TREE_TYPE (s);
s = TYPE_SIZE_UNIT (s);
}
else
s = OMP_CLAUSE_SIZE (c);
if (s == NULL_TREE) if (s == NULL_TREE)
s = TYPE_SIZE_UNIT (TREE_TYPE (ovar)); s = TYPE_SIZE_UNIT (TREE_TYPE (ovar));
s = fold_convert (size_type_node, s); s = fold_convert (size_type_node, s);
@ -15330,6 +15367,11 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
tkind_zero = tkind; tkind_zero = tkind;
} }
break; break;
case OMP_CLAUSE_FIRSTPRIVATE:
gcc_checking_assert (is_gimple_omp_oacc (ctx->stmt));
tkind = GOMP_MAP_TO;
tkind_zero = tkind;
break;
case OMP_CLAUSE_TO: case OMP_CLAUSE_TO:
tkind = GOMP_MAP_TO; tkind = GOMP_MAP_TO;
tkind_zero = tkind; tkind_zero = tkind;
@ -15369,6 +15411,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
break; break;
case OMP_CLAUSE_FIRSTPRIVATE: case OMP_CLAUSE_FIRSTPRIVATE:
if (is_oacc_parallel (ctx))
goto oacc_firstprivate_map;
ovar = OMP_CLAUSE_DECL (c); ovar = OMP_CLAUSE_DECL (c);
if (is_reference (ovar)) if (is_reference (ovar))
talign = TYPE_ALIGN_UNIT (TREE_TYPE (TREE_TYPE (ovar))); talign = TYPE_ALIGN_UNIT (TREE_TYPE (TREE_TYPE (ovar)));
@ -15543,6 +15587,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
gimple_seq_add_stmt (&new_body, gimple_seq_add_stmt (&new_body,
gimple_build_assign (ctx->receiver_decl, t)); gimple_build_assign (ctx->receiver_decl, t));
} }
gimple_seq_add_seq (&new_body, fplist);
if (offloaded || data_region) if (offloaded || data_region)
{ {
@ -15554,6 +15599,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
default: default:
break; break;
case OMP_CLAUSE_FIRSTPRIVATE: case OMP_CLAUSE_FIRSTPRIVATE:
if (is_gimple_omp_oacc (ctx->stmt))
break;
var = OMP_CLAUSE_DECL (c); var = OMP_CLAUSE_DECL (c);
if (is_reference (var) if (is_reference (var)
|| is_gimple_reg_type (TREE_TYPE (var))) || is_gimple_reg_type (TREE_TYPE (var)))
@ -15639,6 +15686,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
} }
break; break;
case OMP_CLAUSE_PRIVATE: case OMP_CLAUSE_PRIVATE:
if (is_gimple_omp_oacc (ctx->stmt))
break;
var = OMP_CLAUSE_DECL (c); var = OMP_CLAUSE_DECL (c);
if (is_reference (var)) if (is_reference (var))
{ {
@ -15727,7 +15776,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
/* Handle GOMP_MAP_FIRSTPRIVATE_{POINTER,REFERENCE} in second pass, /* Handle GOMP_MAP_FIRSTPRIVATE_{POINTER,REFERENCE} in second pass,
so that firstprivate vars holding OMP_CLAUSE_SIZE if needed so that firstprivate vars holding OMP_CLAUSE_SIZE if needed
are already handled. */ are already handled. */
for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c)) for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
switch (OMP_CLAUSE_CODE (c)) switch (OMP_CLAUSE_CODE (c))
{ {
tree var; tree var;

5
gcc/testsuite/ChangeLog
View File

@ -1,3 +1,8 @@
2015-11-11 Nathan Sidwell <nathan@codesourcery.com>
* gfortran.dg/goacc/private-3.f95: Remove xfail.
* gfortran.dg/goacc/combined_loop.f90: Remove xfail.
2015-11-11 Eric Botcazou <ebotcazou@adacore.com> 2015-11-11 Eric Botcazou <ebotcazou@adacore.com>
* gcc.target/i386/pr67265.c: New test. * gcc.target/i386/pr67265.c: New test.

2
gcc/testsuite/gfortran.dg/goacc/combined_loop.f90
View File

@ -1,6 +1,4 @@
! { dg-do compile } ! { dg-do compile }
! <http://news.gmane.org/find-root.php?message_id=%3C563B78B5.5090506%40acm.org%3E>
! { dg-xfail-if "TODO" { *-*-* } }
! !
! PR fortran/64726 ! PR fortran/64726

2
gcc/testsuite/gfortran.dg/goacc/private-3.f95
View File

@ -1,6 +1,4 @@
! { dg-do compile } ! { dg-do compile }
! <http://news.gmane.org/find-root.php?message_id=%3C563B78B5.5090506%40acm.org%3E>
! { dg-xfail-if "TODO" { *-*-* } }
! test for private variables in a reduction clause ! test for private variables in a reduction clause

5
libgomp/ChangeLog
View File

@ -1,3 +1,8 @@
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-2.c: New.
2015-11-09 Nathan Sidwell <nathan@codesourcery.com> 2015-11-09 Nathan Sidwell <nathan@codesourcery.com>
* testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c: Remove * testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c: Remove

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

@ -0,0 +1,41 @@
/* { dg-do run } */
#include <openacc.h>
int main ()
{
int ok = 1;
int val = 2;
int ary[32];
int ondev = 0;
for (int i = 0; i < 32; i++)
ary[i] = ~0;
#pragma acc parallel num_gangs (32) copy (ok) firstprivate (val) copy(ary, 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;
}

31
libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-2.c Normal file
View File

@ -0,0 +1,31 @@
/* { dg-do run } */
#include <openacc.h>
int main ()
{
int ok = 1;
int val = 2;
#pragma acc data copy(val)
{
#pragma acc parallel present (val)
{
val = 7;
}
#pragma acc parallel firstprivate (val) copy(ok)
{
ok = val == 7;
val = 9;
}
}
if (!ok)
return 1;
if(val != 7)
return 1;
return 0;
}

2
libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c
View File

@ -1,7 +1,5 @@
/* { dg-do run } */ /* { dg-do run } */
/* { dg-additional-options "-O2" */ /* { dg-additional-options "-O2" */
/* <http://news.gmane.org/find-root.php?message_id=%3C563B78B5.5090506%40acm.org%3E>
{ dg-xfail-if "TODO" { *-*-* } } */
#include <stdio.h> #include <stdio.h>

2
libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c
View File

@ -1,7 +1,5 @@
/* { dg-do run } */ /* { dg-do run } */
/* { dg-additional-options "-O2" */ /* { dg-additional-options "-O2" */
/* <http://news.gmane.org/find-root.php?message_id=%3C563B78B5.5090506%40acm.org%3E>
{ dg-xfail-if "TODO" { *-*-* } } */
#include <stdio.h> #include <stdio.h>