+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>
PR target/67265
enum omp_region_type
{
- ORT_WORKSHARE = 0,
- ORT_SIMD = 1,
- ORT_PARALLEL = 2,
- ORT_COMBINED_PARALLEL = 3,
- ORT_TASK = 4,
- ORT_UNTIED_TASK = 5,
- ORT_TEAMS = 8,
- ORT_COMBINED_TEAMS = 9,
+ ORT_WORKSHARE = 0x00,
+ ORT_SIMD = 0x01,
+
+ ORT_PARALLEL = 0x02,
+ ORT_COMBINED_PARALLEL = 0x03,
+
+ ORT_TASK = 0x04,
+ ORT_UNTIED_TASK = 0x05,
+
+ ORT_TEAMS = 0x08,
+ ORT_COMBINED_TEAMS = 0x09,
+
/* Data region. */
- ORT_TARGET_DATA = 16,
+ ORT_TARGET_DATA = 0x10,
+
/* Data region with offloading. */
- ORT_TARGET = 32,
- ORT_COMBINED_TARGET = 33,
+ ORT_TARGET = 0x20,
+ 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
DECL_VALUE_EXPRs in taskloop pre body. */
- ORT_NONE = 64
+ ORT_NONE = 0x100
};
/* Gimplify hashtable helper. */
struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
while (ctx
&& (ctx->region_type == ORT_WORKSHARE
- || ctx->region_type == ORT_SIMD))
+ || ctx->region_type == ORT_SIMD
+ || ctx->region_type == ORT_ACC))
ctx = ctx->outer_context;
if (ctx)
omp_add_variable (ctx, tmp, GOVD_LOCAL | GOVD_SEEN);
struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
while (ctx
&& (ctx->region_type == ORT_WORKSHARE
- || ctx->region_type == ORT_SIMD))
+ || ctx->region_type == ORT_SIMD
+ || ctx->region_type == ORT_ACC))
ctx = ctx->outer_context;
if (!ctx && !nonlocal_vlas->add (decl))
{
}
else if (ctx->region_type != ORT_WORKSHARE
&& 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);
ctx = ctx->outer_context;
/* We shouldn't be re-adding the decl with the same data
sharing class. */
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;
- gcc_assert ((nflags & GOVD_DATA_SHARE_CLASS)
- == (GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE)
+ /* The only combination of data sharing classes we should see is
+ 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);
n->value = nflags;
return;
else if (is_scalar)
nflags |= GOVD_FIRSTPRIVATE;
}
- 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))
+
+ struct gimplify_omp_ctx *octx = ctx->outer_context;
+ if ((ctx->region_type & ORT_ACC) && octx)
{
- error ("%qD referenced in target region does not have "
- "a mappable type", decl);
- nflags |= GOVD_MAP | GOVD_EXPLICIT;
+ /* Look in outer OpenACC contexts, to see if there's a
+ data attribute for this variable. */
+ 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);
}
else
{
if (ctx->region_type == ORT_WORKSHARE
|| 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;
flags = omp_default_clause (ctx, decl, in_code, flags);
}
if (ctx->region_type != ORT_WORKSHARE
- && ctx->region_type != ORT_SIMD)
+ && ctx->region_type != ORT_SIMD
+ && ctx->region_type != ORT_ACC)
return false;
else if (ctx->outer_context)
return omp_is_private (ctx->outer_context, decl, simd);
}
}
while (ctx->region_type == ORT_WORKSHARE
- || ctx->region_type == ORT_SIMD);
+ || ctx->region_type == ORT_SIMD
+ || ctx->region_type == ORT_ACC);
return false;
}
omp_notice_variable (outer_ctx->outer_context, decl, true);
}
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
&& splay_tree_lookup (outer_ctx->variables,
(splay_tree_key) decl) == NULL
goto do_add;
case OMP_CLAUSE_REDUCTION:
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);
if (TREE_CODE (decl) == MEM_REF)
{
{
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);
gimplify_adjust_omp_clauses (pre_p, &OACC_CACHE_CLAUSES (expr), OACC_CACHE);
case OMP_FOR:
case CILK_FOR:
case OMP_DISTRIBUTE:
+ break;
case OACC_LOOP:
+ ort = ORT_ACC;
break;
case OMP_TASKLOOP:
if (find_omp_clause (OMP_FOR_CLAUSES (for_stmt), OMP_CLAUSE_UNTIED))
ort = OMP_TARGET_COMBINED (expr) ? ORT_COMBINED_TARGET : ORT_TARGET;
break;
case OACC_KERNELS:
+ ort = ORT_ACC_KERNELS;
+ break;
case OACC_PARALLEL:
- ort = ORT_TARGET;
+ ort = ORT_ACC_PARALLEL;
break;
case OACC_DATA:
+ ort = ORT_ACC_DATA;
+ break;
case OMP_TARGET_DATA:
ort = ORT_TARGET_DATA;
break;
pop_gimplify_context (g);
else
pop_gimplify_context (NULL);
- if (ort == ORT_TARGET_DATA)
+ if ((ort & ORT_TARGET_DATA) != 0)
{
enum built_in_function end_ix;
switch (TREE_CODE (expr))
tree expr = *expr_p;
int kind;
gomp_target *stmt;
+ enum omp_region_type ort = ORT_WORKSHARE;
switch (TREE_CODE (expr))
{
case OACC_ENTER_DATA:
- kind = GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA;
- break;
case OACC_EXIT_DATA:
kind = GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA;
+ ort = ORT_ACC;
break;
case OACC_UPDATE:
kind = GF_OMP_TARGET_KIND_OACC_UPDATE;
+ ort = ORT_ACC;
break;
case OMP_TARGET_UPDATE:
kind = GF_OMP_TARGET_KIND_UPDATE;
gcc_unreachable ();
}
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),
TREE_CODE (expr));
stmt = gimple_build_omp_target (NULL, kind, OMP_STANDALONE_CLAUSES (expr));
/* FALLTHRU */
case OMP_CLAUSE_FIRSTPRIVATE:
- if (is_gimple_omp_oacc (ctx->stmt))
- {
- sorry ("clause not supported yet");
- break;
- }
- /* FALLTHRU */
case OMP_CLAUSE_LINEAR:
decl = OMP_CLAUSE_DECL (c);
do_private:
/* FALLTHRU */
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_LINEAR:
case OMP_CLAUSE_IS_DEVICE_PTR:
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,
allocate new backing storage for the new pointer
tree child_fn, t, c;
gomp_target *stmt = as_a <gomp_target *> (gsi_stmt (*gsi_p));
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);
bool offloaded, data_region;
unsigned int map_cnt = 0;
child_fn = ctx->cb.dst_fn;
push_gimplify_context ();
+ fplist = NULL;
for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
switch (OMP_CLAUSE_CODE (c))
/* FALLTHRU */
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
+ oacc_firstprivate:
var = OMP_CLAUSE_DECL (c);
if (!DECL_P (var))
{
}
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_REFERENCE))
{
x = build_receiver_ref (var, true, 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)
&& TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
x = build_simple_mem_ref (x);
- SET_DECL_VALUE_EXPR (new_var, x);
- DECL_HAS_VALUE_EXPR_P (new_var) = 1;
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
+ {
+ 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++;
break;
case OMP_CLAUSE_FIRSTPRIVATE:
+ if (is_oacc_parallel (ctx))
+ goto oacc_firstprivate;
map_cnt++;
var = OMP_CLAUSE_DECL (c);
if (!is_reference (var)
break;
case OMP_CLAUSE_PRIVATE:
+ if (is_gimple_omp_oacc (ctx->stmt))
+ break;
var = OMP_CLAUSE_DECL (c);
if (is_variable_sized (var))
{
default:
break;
+
case OMP_CLAUSE_MAP:
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
+ oacc_firstprivate_map:
nc = c;
ovar = OMP_CLAUSE_DECL (c);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
x = build_sender_ref (ovar, ctx);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
- && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
- && TREE_CODE (TREE_TYPE (ovar)) == ARRAY_TYPE)
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
+ && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
+ && TREE_CODE (TREE_TYPE (ovar)) == ARRAY_TYPE)
{
gcc_assert (offloaded);
tree avar
avar = build_fold_addr_expr (avar);
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))
{
gcc_assert (offloaded);
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)
s = TYPE_SIZE_UNIT (TREE_TYPE (ovar));
s = fold_convert (size_type_node, s);
tkind_zero = tkind;
}
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:
tkind = GOMP_MAP_TO;
tkind_zero = tkind;
break;
case OMP_CLAUSE_FIRSTPRIVATE:
+ if (is_oacc_parallel (ctx))
+ goto oacc_firstprivate_map;
ovar = OMP_CLAUSE_DECL (c);
if (is_reference (ovar))
talign = TYPE_ALIGN_UNIT (TREE_TYPE (TREE_TYPE (ovar)));
gimple_seq_add_stmt (&new_body,
gimple_build_assign (ctx->receiver_decl, t));
}
+ gimple_seq_add_seq (&new_body, fplist);
if (offloaded || data_region)
{
default:
break;
case OMP_CLAUSE_FIRSTPRIVATE:
+ if (is_gimple_omp_oacc (ctx->stmt))
+ break;
var = OMP_CLAUSE_DECL (c);
if (is_reference (var)
|| is_gimple_reg_type (TREE_TYPE (var)))
}
break;
case OMP_CLAUSE_PRIVATE:
+ if (is_gimple_omp_oacc (ctx->stmt))
+ break;
var = OMP_CLAUSE_DECL (c);
if (is_reference (var))
{
/* Handle GOMP_MAP_FIRSTPRIVATE_{POINTER,REFERENCE} in second pass,
so that firstprivate vars holding OMP_CLAUSE_SIZE if needed
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))
{
tree var;
+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>
* gcc.target/i386/pr67265.c: New test.
! { dg-do compile }
-! <http://news.gmane.org/find-root.php?message_id=%3C563B78B5.5090506%40acm.org%3E>
-! { dg-xfail-if "TODO" { *-*-* } }
!
! PR fortran/64726
! { 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
+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>
* testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c: Remove
--- /dev/null
+/* { 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;
+}
--- /dev/null
+/* { 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;
+}
/* { dg-do run } */
/* { 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>
/* { dg-do run } */
/* { 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>
projects / gcc.git / commitdiff