[gomp4.5] Support lastprivate on distribute
Jakub Jelinek
jakub@redhat.com
Fri Oct 23 12:16:00 GMT 2015
Hi!
This patch adds support for lastprivate clause on distribute
construct (firstprivate+lastprivate on the same var
is disallowed, as one can't synchronize between contention groups).
Regtested on x86_64-linux (normal and intelmicemul offloading).
2015-10-23 Jakub Jelinek <jakub@redhat.com>
gcc/
* gimplify.c (omp_default_clause): Tweak for
* private/firstprivate/is_device_ptr variables on target
construct and use_device_ptr on target data.
(omp_check_private): Likewise.
(omp_no_lastprivate): Return true only for Fortran.
(gimplify_scan_omp_clauses): Fix up handling of lastprivate
and linear when combined with distribute.
(gimplify_adjust_omp_clauses): Diagnose the same var
on both firstprivate and lastprivate on distribute construct.
(gimplify_omp_for): Fix up handling of predetermined
lastprivate or linear iter vars when combined with distribute.
* omp-low.c (add_taskreg_looptemp_clauses): Add one extra
_looptemp_ clause even for distribute parallel for, if there
are lastprivate clauses on the for.
(expand_omp_for_static_nochunk, expand_omp_for_static_chunk):
Initialize the extra _looptemp_ clause to fd->loop.n2.
(lower_omp_for_lastprivate): Determine the right count variable
for distribute simd, or distribute parallel for{, simd}.
gcc/c-family/
* c-omp.c (c_omp_split_clauses): Adjust for lastprivate being
allowed on distribute.
gcc/c/
* c-parser.c (OMP_DISTRIBUTE_CLAUSE_MASK): Add lastprivate clause.
gcc/cp/
* parser.c (OMP_DISTRIBUTE_CLAUSE_MASK): Add lastprivate clause.
gcc/testsuite/
* c-c++-common/gomp/distribute-1.c: New test.
* c-c++-common/gomp/pr61486-2.c: Add #pragma omp declare target
and #pragma omp end declare target pair around the function.
Change s from a parameter to a file scope variable.
libgomp/
* testsuite/libgomp.c/pr66199-5.c: New test.
* testsuite/libgomp.c/pr66199-6.c: New test.
* testsuite/libgomp.c/pr66199-7.c: New test.
* testsuite/libgomp.c/pr66199-8.c: New test.
* testsuite/libgomp.c/pr66199-9.c: New test.
* testsuite/libgomp.c++/pr66199-3.C: New test.
* testsuite/libgomp.c++/pr66199-4.C: New test.
* testsuite/libgomp.c++/pr66199-5.C: New test.
* testsuite/libgomp.c++/pr66199-6.C: New test.
* testsuite/libgomp.c++/pr66199-7.C: New test.
* testsuite/libgomp.c++/pr66199-8.C: New test.
* testsuite/libgomp.c++/pr66199-9.C: New test.
--- gcc/gimplify.c.jj 2015-10-22 13:41:23.214882261 +0200
+++ gcc/gimplify.c 2015-10-22 18:01:45.125588093 +0200
@@ -5849,9 +5849,10 @@ omp_default_clause (struct gimplify_omp_
{
splay_tree_node n2;
- if ((octx->region_type & (ORT_TARGET_DATA | ORT_TARGET)) != 0)
- continue;
n2 = splay_tree_lookup (octx->variables, (splay_tree_key) decl);
+ if ((octx->region_type & (ORT_TARGET_DATA | ORT_TARGET)) != 0
+ && (n2 == NULL || (n2->value & GOVD_DATA_SHARE_CLASS) == 0))
+ continue;
if (n2 && (n2->value & GOVD_DATA_SHARE_CLASS) != GOVD_SHARED)
{
flags |= GOVD_FIRSTPRIVATE;
@@ -6143,10 +6144,12 @@ omp_check_private (struct gimplify_omp_c
return true;
}
- if ((ctx->region_type & (ORT_TARGET | ORT_TARGET_DATA)) != 0)
+ n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
+
+ if ((ctx->region_type & (ORT_TARGET | ORT_TARGET_DATA)) != 0
+ && (n == NULL || (n->value & GOVD_DATA_SHARE_CLASS) == 0))
continue;
- n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
if (n != NULL)
{
if ((n->value & GOVD_LOCAL) != 0
@@ -6177,12 +6180,12 @@ omp_no_lastprivate (struct gimplify_omp_
if (!ctx->combined_loop)
return false;
if (ctx->distribute)
- return true;
+ return lang_GNU_Fortran ();
break;
case ORT_COMBINED_PARALLEL:
break;
case ORT_COMBINED_TEAMS:
- return true;
+ return lang_GNU_Fortran ();
default:
return false;
}
@@ -6279,16 +6282,25 @@ gimplify_scan_omp_clauses (tree *list_p,
else if (error_operand_p (decl))
goto do_add;
else if (outer_ctx
- && outer_ctx->region_type == ORT_COMBINED_PARALLEL
+ && (outer_ctx->region_type == ORT_COMBINED_PARALLEL
+ || outer_ctx->region_type == ORT_COMBINED_TEAMS)
&& splay_tree_lookup (outer_ctx->variables,
(splay_tree_key) decl) == NULL)
- omp_add_variable (outer_ctx, decl, GOVD_SHARED | GOVD_SEEN);
+ {
+ omp_add_variable (outer_ctx, decl, GOVD_SHARED | GOVD_SEEN);
+ if (outer_ctx->outer_context)
+ omp_notice_variable (outer_ctx->outer_context, decl, true);
+ }
else if (outer_ctx
&& (outer_ctx->region_type & ORT_TASK) != 0
&& outer_ctx->combined_loop
&& splay_tree_lookup (outer_ctx->variables,
(splay_tree_key) decl) == NULL)
- omp_add_variable (outer_ctx, decl, GOVD_LASTPRIVATE | GOVD_SEEN);
+ {
+ omp_add_variable (outer_ctx, decl, GOVD_LASTPRIVATE | GOVD_SEEN);
+ if (outer_ctx->outer_context)
+ omp_notice_variable (outer_ctx->outer_context, decl, true);
+ }
else if (outer_ctx
&& outer_ctx->region_type == ORT_WORKSHARE
&& outer_ctx->combined_loop
@@ -6302,8 +6314,14 @@ gimplify_scan_omp_clauses (tree *list_p,
== ORT_COMBINED_PARALLEL)
&& splay_tree_lookup (outer_ctx->outer_context->variables,
(splay_tree_key) decl) == NULL)
- omp_add_variable (outer_ctx->outer_context, decl,
- GOVD_SHARED | GOVD_SEEN);
+ {
+ struct gimplify_omp_ctx *octx = outer_ctx->outer_context;
+ omp_add_variable (octx, decl, GOVD_SHARED | GOVD_SEEN);
+ if (octx->outer_context)
+ omp_notice_variable (octx->outer_context, decl, true);
+ }
+ else if (outer_ctx->outer_context)
+ omp_notice_variable (outer_ctx->outer_context, decl, true);
}
goto do_add;
case OMP_CLAUSE_REDUCTION:
@@ -6416,9 +6434,7 @@ gimplify_scan_omp_clauses (tree *list_p,
{
if (octx->outer_context
&& (octx->outer_context->region_type
- == ORT_COMBINED_PARALLEL
- || (octx->outer_context->region_type
- == ORT_COMBINED_TEAMS)))
+ == ORT_COMBINED_PARALLEL))
octx = octx->outer_context;
else if (omp_check_private (octx, decl, false))
break;
@@ -6433,8 +6449,15 @@ gimplify_scan_omp_clauses (tree *list_p,
&& octx == outer_ctx)
flags = GOVD_SEEN | GOVD_SHARED;
else if (octx
+ && octx->region_type == ORT_COMBINED_TEAMS)
+ flags = GOVD_SEEN | GOVD_SHARED;
+ else if (octx
&& octx->region_type == ORT_COMBINED_TARGET)
- flags &= ~GOVD_LASTPRIVATE;
+ {
+ flags &= ~GOVD_LASTPRIVATE;
+ if (flags == GOVD_SEEN)
+ break;
+ }
else
break;
splay_tree_node on
@@ -7289,6 +7312,15 @@ gimplify_adjust_omp_clauses (gimple_seq
else
OMP_CLAUSE_CODE (c) = OMP_CLAUSE_PRIVATE;
}
+ else if (code == OMP_DISTRIBUTE
+ && OMP_CLAUSE_LASTPRIVATE_FIRSTPRIVATE (c))
+ {
+ remove = true;
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "same variable used in %<firstprivate%> and "
+ "%<lastprivate%> clauses on %<distribute%> "
+ "construct");
+ }
break;
case OMP_CLAUSE_ALIGNED:
@@ -7902,6 +7934,26 @@ gimplify_omp_for (tree *expr_p, gimple_s
OMP_CLAUSE_LINEAR_NO_COPYOUT (c) = 1;
flags |= GOVD_LINEAR_LASTPRIVATE_NO_OUTER;
}
+ else
+ {
+ struct gimplify_omp_ctx *octx = outer->outer_context;
+ if (octx
+ && octx->region_type == ORT_COMBINED_PARALLEL
+ && octx->outer_context
+ && (octx->outer_context->region_type
+ == ORT_WORKSHARE)
+ && octx->outer_context->combined_loop)
+ {
+ octx = octx->outer_context;
+ n = splay_tree_lookup (octx->variables,
+ (splay_tree_key)decl);
+ if (n != NULL && (n->value & GOVD_LOCAL) != 0)
+ {
+ OMP_CLAUSE_LINEAR_NO_COPYOUT (c) = 1;
+ flags |= GOVD_LINEAR_LASTPRIVATE_NO_OUTER;
+ }
+ }
+ }
}
}
@@ -7936,7 +7988,41 @@ gimplify_omp_for (tree *expr_p, gimple_s
{
omp_add_variable (outer, decl,
GOVD_LASTPRIVATE | GOVD_SEEN);
- if (outer->outer_context)
+ if (outer->region_type == ORT_COMBINED_PARALLEL
+ && outer->outer_context
+ && (outer->outer_context->region_type
+ == ORT_WORKSHARE)
+ && outer->outer_context->combined_loop)
+ {
+ outer = outer->outer_context;
+ n = splay_tree_lookup (outer->variables,
+ (splay_tree_key)decl);
+ if (omp_check_private (outer, decl, false))
+ outer = NULL;
+ else if (n == NULL
+ || ((n->value & GOVD_DATA_SHARE_CLASS)
+ == 0))
+ omp_add_variable (outer, decl,
+ GOVD_LASTPRIVATE
+ | GOVD_SEEN);
+ else
+ outer = NULL;
+ }
+ if (outer && outer->outer_context
+ && (outer->outer_context->region_type
+ == ORT_COMBINED_TEAMS))
+ {
+ outer = outer->outer_context;
+ n = splay_tree_lookup (outer->variables,
+ (splay_tree_key)decl);
+ if (n == NULL
+ || (n->value & GOVD_DATA_SHARE_CLASS) == 0)
+ omp_add_variable (outer, decl,
+ GOVD_SHARED | GOVD_SEEN);
+ else
+ outer = NULL;
+ }
+ if (outer && outer->outer_context)
omp_notice_variable (outer->outer_context, decl,
true);
}
@@ -7985,7 +8071,41 @@ gimplify_omp_for (tree *expr_p, gimple_s
{
omp_add_variable (outer, decl,
GOVD_LASTPRIVATE | GOVD_SEEN);
- if (outer->outer_context)
+ if (outer->region_type == ORT_COMBINED_PARALLEL
+ && outer->outer_context
+ && (outer->outer_context->region_type
+ == ORT_WORKSHARE)
+ && outer->outer_context->combined_loop)
+ {
+ outer = outer->outer_context;
+ n = splay_tree_lookup (outer->variables,
+ (splay_tree_key)decl);
+ if (omp_check_private (outer, decl, false))
+ outer = NULL;
+ else if (n == NULL
+ || ((n->value & GOVD_DATA_SHARE_CLASS)
+ == 0))
+ omp_add_variable (outer, decl,
+ GOVD_LASTPRIVATE
+ | GOVD_SEEN);
+ else
+ outer = NULL;
+ }
+ if (outer && outer->outer_context
+ && (outer->outer_context->region_type
+ == ORT_COMBINED_TEAMS))
+ {
+ outer = outer->outer_context;
+ n = splay_tree_lookup (outer->variables,
+ (splay_tree_key)decl);
+ if (n == NULL
+ || (n->value & GOVD_DATA_SHARE_CLASS) == 0)
+ omp_add_variable (outer, decl,
+ GOVD_SHARED | GOVD_SEEN);
+ else
+ outer = NULL;
+ }
+ if (outer && outer->outer_context)
omp_notice_variable (outer->outer_context, decl,
true);
}
--- gcc/omp-low.c.jj 2015-10-22 14:11:23.488003543 +0200
+++ gcc/omp-low.c 2015-10-23 13:10:16.995777473 +0200
@@ -2646,12 +2646,15 @@ add_taskreg_looptemp_clauses (enum gf_ma
&& TREE_CODE (fd.loop.n2) != INTEGER_CST)
{
count += fd.collapse - 1;
- /* For taskloop, if there are lastprivate clauses on the inner
+ /* If there are lastprivate clauses on the inner
GIMPLE_OMP_FOR, add one more temporaries for the total number
of iterations (product of count1 ... countN-1). */
- if (msk == GF_OMP_FOR_KIND_TASKLOOP
- && find_omp_clause (gimple_omp_for_clauses (for_stmt),
- OMP_CLAUSE_LASTPRIVATE))
+ if (find_omp_clause (gimple_omp_for_clauses (for_stmt),
+ OMP_CLAUSE_LASTPRIVATE))
+ count++;
+ else if (msk == GF_OMP_FOR_KIND_FOR
+ && find_omp_clause (gimple_omp_parallel_clauses (stmt),
+ OMP_CLAUSE_LASTPRIVATE))
count++;
}
for (i = 0; i < count; i++)
@@ -8648,6 +8651,30 @@ expand_omp_for_static_nochunk (struct om
OMP_CLAUSE__LOOPTEMP_);
gcc_assert (innerc);
endvar = OMP_CLAUSE_DECL (innerc);
+ if (fd->collapse > 1 && TREE_CODE (fd->loop.n2) != INTEGER_CST
+ && gimple_omp_for_kind (fd->for_stmt) == GF_OMP_FOR_KIND_DISTRIBUTE)
+ {
+ int i;
+ for (i = 1; i < fd->collapse; i++)
+ {
+ innerc = find_omp_clause (OMP_CLAUSE_CHAIN (innerc),
+ OMP_CLAUSE__LOOPTEMP_);
+ gcc_assert (innerc);
+ }
+ innerc = find_omp_clause (OMP_CLAUSE_CHAIN (innerc),
+ OMP_CLAUSE__LOOPTEMP_);
+ if (innerc)
+ {
+ /* If needed (distribute parallel for with lastprivate),
+ propagate down the total number of iterations. */
+ tree t = fold_convert (TREE_TYPE (OMP_CLAUSE_DECL (innerc)),
+ fd->loop.n2);
+ t = force_gimple_operand_gsi (&gsi, t, false, NULL_TREE, false,
+ GSI_CONTINUE_LINKING);
+ assign_stmt = gimple_build_assign (OMP_CLAUSE_DECL (innerc), t);
+ gsi_insert_after (&gsi, assign_stmt, GSI_CONTINUE_LINKING);
+ }
+ }
}
t = fold_convert (itype, s0);
t = fold_build2 (MULT_EXPR, itype, t, step);
@@ -9133,6 +9160,30 @@ expand_omp_for_static_chunk (struct omp_
OMP_CLAUSE__LOOPTEMP_);
gcc_assert (innerc);
endvar = OMP_CLAUSE_DECL (innerc);
+ if (fd->collapse > 1 && TREE_CODE (fd->loop.n2) != INTEGER_CST
+ && gimple_omp_for_kind (fd->for_stmt) == GF_OMP_FOR_KIND_DISTRIBUTE)
+ {
+ int i;
+ for (i = 1; i < fd->collapse; i++)
+ {
+ innerc = find_omp_clause (OMP_CLAUSE_CHAIN (innerc),
+ OMP_CLAUSE__LOOPTEMP_);
+ gcc_assert (innerc);
+ }
+ innerc = find_omp_clause (OMP_CLAUSE_CHAIN (innerc),
+ OMP_CLAUSE__LOOPTEMP_);
+ if (innerc)
+ {
+ /* If needed (distribute parallel for with lastprivate),
+ propagate down the total number of iterations. */
+ tree t = fold_convert (TREE_TYPE (OMP_CLAUSE_DECL (innerc)),
+ fd->loop.n2);
+ t = force_gimple_operand_gsi (&gsi, t, false, NULL_TREE, false,
+ GSI_CONTINUE_LINKING);
+ assign_stmt = gimple_build_assign (OMP_CLAUSE_DECL (innerc), t);
+ gsi_insert_after (&gsi, assign_stmt, GSI_CONTINUE_LINKING);
+ }
+ }
}
t = fold_convert (itype, s0);
@@ -13456,26 +13507,36 @@ lower_omp_for_lastprivate (struct omp_fo
&& TREE_CODE (n2) != INTEGER_CST
&& gimple_omp_for_combined_into_p (fd->for_stmt))
{
- struct omp_context *task_ctx = NULL;
+ struct omp_context *taskreg_ctx = NULL;
if (gimple_code (ctx->outer->stmt) == GIMPLE_OMP_FOR)
{
gomp_for *gfor = as_a <gomp_for *> (ctx->outer->stmt);
- if (gimple_omp_for_kind (gfor) == GF_OMP_FOR_KIND_FOR)
+ if (gimple_omp_for_kind (gfor) == GF_OMP_FOR_KIND_FOR
+ || gimple_omp_for_kind (gfor) == GF_OMP_FOR_KIND_DISTRIBUTE)
{
- struct omp_for_data outer_fd;
- extract_omp_for_data (gfor, &outer_fd, NULL);
- n2 = fold_convert (TREE_TYPE (n2), outer_fd.loop.n2);
+ if (gimple_omp_for_combined_into_p (gfor))
+ {
+ gcc_assert (ctx->outer->outer
+ && is_parallel_ctx (ctx->outer->outer));
+ taskreg_ctx = ctx->outer->outer;
+ }
+ else
+ {
+ struct omp_for_data outer_fd;
+ extract_omp_for_data (gfor, &outer_fd, NULL);
+ n2 = fold_convert (TREE_TYPE (n2), outer_fd.loop.n2);
+ }
}
else if (gimple_omp_for_kind (gfor) == GF_OMP_FOR_KIND_TASKLOOP)
- task_ctx = ctx->outer->outer;
+ taskreg_ctx = ctx->outer->outer;
}
- else if (is_task_ctx (ctx->outer))
- task_ctx = ctx->outer;
- if (task_ctx)
+ else if (is_taskreg_ctx (ctx->outer))
+ taskreg_ctx = ctx->outer;
+ if (taskreg_ctx)
{
int i;
tree innerc
- = find_omp_clause (gimple_omp_task_clauses (task_ctx->stmt),
+ = find_omp_clause (gimple_omp_taskreg_clauses (taskreg_ctx->stmt),
OMP_CLAUSE__LOOPTEMP_);
gcc_assert (innerc);
for (i = 0; i < fd->collapse; i++)
@@ -13489,7 +13550,7 @@ lower_omp_for_lastprivate (struct omp_fo
if (innerc)
n2 = fold_convert (TREE_TYPE (n2),
lookup_decl (OMP_CLAUSE_DECL (innerc),
- task_ctx));
+ taskreg_ctx));
}
}
cond = build2 (cond_code, boolean_type_node, fd->loop.v, n2);
--- gcc/c-family/c-omp.c.jj 2015-10-22 13:41:23.256881650 +0200
+++ gcc/c-family/c-omp.c 2015-10-22 14:35:41.075032703 +0200
@@ -1097,10 +1097,24 @@ c_omp_split_clauses (location_t loc, enu
s = C_OMP_CLAUSE_SPLIT_FOR;
}
break;
- /* Lastprivate is allowed on for, sections and simd. In
+ /* Lastprivate is allowed on distribute, for, sections and simd. In
parallel {for{, simd},sections} we actually want to put it on
parallel rather than for or sections. */
case OMP_CLAUSE_LASTPRIVATE:
+ if (code == OMP_DISTRIBUTE)
+ {
+ s = C_OMP_CLAUSE_SPLIT_DISTRIBUTE;
+ break;
+ }
+ if ((mask & (OMP_CLAUSE_MASK_1
+ << PRAGMA_OMP_CLAUSE_DIST_SCHEDULE)) != 0)
+ {
+ c = build_omp_clause (OMP_CLAUSE_LOCATION (clauses),
+ OMP_CLAUSE_LASTPRIVATE);
+ OMP_CLAUSE_DECL (c) = OMP_CLAUSE_DECL (clauses);
+ OMP_CLAUSE_CHAIN (c) = cclauses[C_OMP_CLAUSE_SPLIT_DISTRIBUTE];
+ cclauses[C_OMP_CLAUSE_SPLIT_DISTRIBUTE] = c;
+ }
if (code == OMP_FOR || code == OMP_SECTIONS)
{
if ((mask & (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NUM_THREADS))
--- gcc/c/c-parser.c.jj 2015-10-22 13:41:23.230882028 +0200
+++ gcc/c/c-parser.c 2015-10-22 14:35:41.075032703 +0200
@@ -14682,6 +14682,7 @@ c_parser_omp_cancellation_point (c_parse
#define OMP_DISTRIBUTE_CLAUSE_MASK \
( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_FIRSTPRIVATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_LASTPRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DIST_SCHEDULE)\
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COLLAPSE))
--- gcc/cp/parser.c.jj 2015-10-22 13:41:23.248881766 +0200
+++ gcc/cp/parser.c 2015-10-22 14:35:41.067032818 +0200
@@ -33572,6 +33572,7 @@ cp_parser_omp_cancellation_point (cp_par
#define OMP_DISTRIBUTE_CLAUSE_MASK \
( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_FIRSTPRIVATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_LASTPRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DIST_SCHEDULE)\
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COLLAPSE))
--- gcc/testsuite/c-c++-common/gomp/distribute-1.c.jj 2015-10-22 14:35:41.068032804 +0200
+++ gcc/testsuite/c-c++-common/gomp/distribute-1.c 2015-10-22 14:35:41.068032804 +0200
@@ -0,0 +1,56 @@
+int s1, s2, s3, s4, s5, s6, s7, s8;
+#pragma omp declare target (s1, s2, s3, s4, s5, s6, s7, s8)
+
+void
+f1 (void)
+{
+ int i;
+ #pragma omp distribute
+ for (i = 0; i < 64; i++)
+ ;
+ #pragma omp distribute private (i)
+ for (i = 0; i < 64; i++)
+ ;
+ #pragma omp distribute
+ for (int j = 0; j < 64; j++)
+ ;
+ #pragma omp distribute lastprivate (s1)
+ for (s1 = 0; s1 < 64; s1 += 2)
+ ;
+ #pragma omp distribute lastprivate (s2)
+ for (i = 0; i < 64; i++)
+ s2 = 2 * i;
+ #pragma omp distribute simd
+ for (i = 0; i < 64; i++)
+ ;
+ #pragma omp distribute simd lastprivate (s3, s4) collapse(2)
+ for (s3 = 0; s3 < 64; s3++)
+ for (s4 = 0; s4 < 3; s4++)
+ ;
+ #pragma omp distribute parallel for
+ for (i = 0; i < 64; i++)
+ ;
+ #pragma omp distribute parallel for private (i)
+ for (i = 0; i < 64; i++)
+ ;
+ #pragma omp distribute parallel for lastprivate (s5)
+ for (s5 = 0; s5 < 64; s5++)
+ ;
+ #pragma omp distribute firstprivate (s7) private (s8)
+ for (i = 0; i < 64; i++)
+ s8 = s7++;
+}
+
+void
+f2 (void)
+{
+ int i;
+ #pragma omp distribute lastprivate (i) /* { dg-error "lastprivate variable .i. is private in outer context" } */
+ for (i = 0; i < 64; i++)
+ ;
+ #pragma omp distribute firstprivate (s6) lastprivate (s6) /* { dg-error "same variable used in .firstprivate. and .lastprivate. clauses on .distribute. construct" } */
+ for (i = 0; i < 64; i++)
+ s6 += i;
+}
+
+#pragma omp declare target to(f1, f2)
--- gcc/testsuite/c-c++-common/gomp/pr61486-2.c.jj 2015-10-22 13:41:22.983885625 +0200
+++ gcc/testsuite/c-c++-common/gomp/pr61486-2.c 2015-10-22 14:35:41.068032804 +0200
@@ -355,8 +355,11 @@ test (int n, int o, int p, int q, int r,
int q, i, j;
+#pragma omp declare target
+int s;
+
void
-test2 (int n, int o, int p, int r, int s, int *pp)
+test2 (int n, int o, int p, int r, int *pp)
{
int a[o];
#pragma omp distribute collapse (2) dist_schedule (static, 4) firstprivate (q)
@@ -449,3 +452,4 @@ test2 (int n, int o, int p, int r, int s
s = i * 10;
}
}
+#pragma omp end declare target
--- libgomp/testsuite/libgomp.c/pr66199-5.c.jj 2015-10-22 14:35:41.080032631 +0200
+++ libgomp/testsuite/libgomp.c/pr66199-5.c 2015-10-22 16:15:11.064658921 +0200
@@ -0,0 +1,66 @@
+/* PR middle-end/66199 */
+/* { dg-do run } */
+
+#pragma omp declare target
+int u[1024], v[1024], w[1024];
+#pragma omp end declare target
+
+__attribute__((noinline, noclone)) long
+f1 (long a, long b)
+{
+ long d;
+ #pragma omp target map(from: d)
+ #pragma omp teams distribute parallel for simd default(none) firstprivate (a, b) shared(u, v, w)
+ for (d = a; d < b; d++)
+ u[d] = v[d] + w[d];
+ return d;
+}
+
+__attribute__((noinline, noclone)) long
+f2 (long a, long b, long c)
+{
+ long d, e;
+ #pragma omp target map(from: d, e)
+ #pragma omp teams distribute parallel for simd default(none) firstprivate (a, b, c) shared(u, v, w) linear(d) lastprivate(e)
+ for (d = a; d < b; d++)
+ {
+ u[d] = v[d] + w[d];
+ e = c + d * 5;
+ }
+ return d + e;
+}
+
+__attribute__((noinline, noclone)) long
+f3 (long a1, long b1, long a2, long b2)
+{
+ long d1, d2;
+ #pragma omp target map(from: d1, d2)
+ #pragma omp teams distribute parallel for simd default(none) firstprivate (a1, b1, a2, b2) shared(u, v, w) lastprivate(d1, d2) collapse(2)
+ for (d1 = a1; d1 < b1; d1++)
+ for (d2 = a2; d2 < b2; d2++)
+ u[d1 * 32 + d2] = v[d1 * 32 + d2] + w[d1 * 32 + d2];
+ return d1 + d2;
+}
+
+__attribute__((noinline, noclone)) long
+f4 (long a1, long b1, long a2, long b2)
+{
+ long d1, d2;
+ #pragma omp target map(from: d1, d2)
+ #pragma omp teams distribute parallel for simd default(none) firstprivate (a1, b1, a2, b2) shared(u, v, w) collapse(2)
+ for (d1 = a1; d1 < b1; d1++)
+ for (d2 = a2; d2 < b2; d2++)
+ u[d1 * 32 + d2] = v[d1 * 32 + d2] + w[d1 * 32 + d2];
+ return d1 + d2;
+}
+
+int
+main ()
+{
+ if (f1 (0, 1024) != 1024
+ || f2 (0, 1024, 17) != 1024 + (17 + 5 * 1023)
+ || f3 (0, 32, 0, 32) != 64
+ || f4 (0, 32, 0, 32) != 64)
+ __builtin_abort ();
+ return 0;
+}
--- libgomp/testsuite/libgomp.c/pr66199-6.c.jj 2015-10-22 14:35:41.080032631 +0200
+++ libgomp/testsuite/libgomp.c/pr66199-6.c 2015-10-22 16:24:37.000000000 +0200
@@ -0,0 +1,42 @@
+/* PR middle-end/66199 */
+/* { dg-do run } */
+/* { dg-options "-O2 -fopenmp" } */
+
+#pragma omp declare target
+int u[1024], v[1024], w[1024];
+#pragma omp end declare target
+
+__attribute__((noinline, noclone)) long
+f2 (long a, long b, long c)
+{
+ long d, e;
+ #pragma omp target map(from: d, e)
+ #pragma omp teams distribute parallel for default(none) firstprivate (a, b, c) shared(u, v, w) lastprivate(d, e)
+ for (d = a; d < b; d++)
+ {
+ u[d] = v[d] + w[d];
+ e = c + d * 5;
+ }
+ return d + e;
+}
+
+__attribute__((noinline, noclone)) long
+f3 (long a1, long b1, long a2, long b2)
+{
+ long d1, d2;
+ #pragma omp target map(from: d1, d2)
+ #pragma omp teams distribute parallel for default(none) firstprivate (a1, b1, a2, b2) shared(u, v, w) lastprivate(d1, d2) collapse(2)
+ for (d1 = a1; d1 < b1; d1++)
+ for (d2 = a2; d2 < b2; d2++)
+ u[d1 * 32 + d2] = v[d1 * 32 + d2] + w[d1 * 32 + d2];
+ return d1 + d2;
+}
+
+int
+main ()
+{
+ if (f2 (0, 1024, 17) != 1024 + (17 + 5 * 1023)
+ || f3 (0, 32, 0, 32) != 64)
+ __builtin_abort ();
+ return 0;
+}
--- libgomp/testsuite/libgomp.c/pr66199-7.c.jj 2015-10-22 14:35:41.080032631 +0200
+++ libgomp/testsuite/libgomp.c/pr66199-7.c 2015-10-22 17:11:33.419476232 +0200
@@ -0,0 +1,66 @@
+/* PR middle-end/66199 */
+/* { dg-do run } */
+
+#pragma omp declare target
+int u[1024], v[1024], w[1024];
+#pragma omp end declare target
+
+__attribute__((noinline, noclone)) long
+f1 (long a, long b)
+{
+ long d;
+ #pragma omp target map(from: d)
+ #pragma omp teams distribute simd default(none) firstprivate (a, b) shared(u, v, w)
+ for (d = a; d < b; d++)
+ u[d] = v[d] + w[d];
+ return d;
+}
+
+__attribute__((noinline, noclone)) long
+f2 (long a, long b, long c)
+{
+ long d, e;
+ #pragma omp target map(from: d, e)
+ #pragma omp teams distribute simd default(none) firstprivate (a, b, c) shared(u, v, w) linear(d) lastprivate(e)
+ for (d = a; d < b; d++)
+ {
+ u[d] = v[d] + w[d];
+ e = c + d * 5;
+ }
+ return d + e;
+}
+
+__attribute__((noinline, noclone)) long
+f3 (long a1, long b1, long a2, long b2)
+{
+ long d1, d2;
+ #pragma omp target map(from: d1, d2)
+ #pragma omp teams distribute simd default(none) firstprivate (a1, b1, a2, b2) shared(u, v, w) lastprivate(d1, d2) collapse(2)
+ for (d1 = a1; d1 < b1; d1++)
+ for (d2 = a2; d2 < b2; d2++)
+ u[d1 * 32 + d2] = v[d1 * 32 + d2] + w[d1 * 32 + d2];
+ return d1 + d2;
+}
+
+__attribute__((noinline, noclone)) long
+f4 (long a1, long b1, long a2, long b2)
+{
+ long d1, d2;
+ #pragma omp target map(from: d1, d2)
+ #pragma omp teams distribute simd default(none) firstprivate (a1, b1, a2, b2) shared(u, v, w) collapse(2)
+ for (d1 = a1; d1 < b1; d1++)
+ for (d2 = a2; d2 < b2; d2++)
+ u[d1 * 32 + d2] = v[d1 * 32 + d2] + w[d1 * 32 + d2];
+ return d1 + d2;
+}
+
+int
+main ()
+{
+ if (f1 (0, 1024) != 1024
+ || f2 (0, 1024, 17) != 1024 + (17 + 5 * 1023)
+ || f3 (0, 32, 0, 32) != 64
+ || f4 (0, 32, 0, 32) != 64)
+ __builtin_abort ();
+ return 0;
+}
--- libgomp/testsuite/libgomp.c/pr66199-8.c.jj 2015-10-22 14:35:41.080032631 +0200
+++ libgomp/testsuite/libgomp.c/pr66199-8.c 2015-10-22 17:11:33.419476232 +0200
@@ -0,0 +1,70 @@
+/* PR middle-end/66199 */
+/* { dg-do run } */
+
+#pragma omp declare target
+int u[1024], v[1024], w[1024];
+#pragma omp end declare target
+
+__attribute__((noinline, noclone)) long
+f1 (long a, long b)
+{
+ long d;
+ #pragma omp target map(from: d)
+ #pragma omp teams default(none) shared(a, b, d, u, v, w)
+ #pragma omp distribute simd firstprivate (a, b)
+ for (d = a; d < b; d++)
+ u[d] = v[d] + w[d];
+ return d;
+}
+
+__attribute__((noinline, noclone)) long
+f2 (long a, long b, long c)
+{
+ long d, e;
+ #pragma omp target map(from: d, e)
+ #pragma omp teams default(none) firstprivate (a, b, c) shared(d, e, u, v, w)
+ #pragma omp distribute simd linear(d) lastprivate(e)
+ for (d = a; d < b; d++)
+ {
+ u[d] = v[d] + w[d];
+ e = c + d * 5;
+ }
+ return d + e;
+}
+
+__attribute__((noinline, noclone)) long
+f3 (long a1, long b1, long a2, long b2)
+{
+ long d1, d2;
+ #pragma omp target map(from: d1, d2)
+ #pragma omp teams default(none) shared(a1, b1, a2, b2, d1, d2, u, v, w)
+ #pragma omp distribute simd firstprivate (a1, b1, a2, b2) lastprivate(d1, d2) collapse(2)
+ for (d1 = a1; d1 < b1; d1++)
+ for (d2 = a2; d2 < b2; d2++)
+ u[d1 * 32 + d2] = v[d1 * 32 + d2] + w[d1 * 32 + d2];
+ return d1 + d2;
+}
+
+__attribute__((noinline, noclone)) long
+f4 (long a1, long b1, long a2, long b2)
+{
+ long d1, d2;
+ #pragma omp target map(from: d1, d2)
+ #pragma omp teams default(none) firstprivate (a1, b1, a2, b2) shared(d1, d2, u, v, w)
+ #pragma omp distribute simd collapse(2)
+ for (d1 = a1; d1 < b1; d1++)
+ for (d2 = a2; d2 < b2; d2++)
+ u[d1 * 32 + d2] = v[d1 * 32 + d2] + w[d1 * 32 + d2];
+ return d1 + d2;
+}
+
+int
+main ()
+{
+ if (f1 (0, 1024) != 1024
+ || f2 (0, 1024, 17) != 1024 + (17 + 5 * 1023)
+ || f3 (0, 32, 0, 32) != 64
+ || f4 (0, 32, 0, 32) != 64)
+ __builtin_abort ();
+ return 0;
+}
--- libgomp/testsuite/libgomp.c/pr66199-9.c.jj 2015-10-22 17:11:56.362148829 +0200
+++ libgomp/testsuite/libgomp.c/pr66199-9.c 2015-10-22 17:12:59.597246434 +0200
@@ -0,0 +1,43 @@
+/* PR middle-end/66199 */
+/* { dg-do run } */
+
+#pragma omp declare target
+int u[1024], v[1024], w[1024];
+#pragma omp end declare target
+
+__attribute__((noinline, noclone)) long
+f2 (long a, long b, long c)
+{
+ long d, e;
+ #pragma omp target map(from: d, e)
+ #pragma omp teams default(none) firstprivate (a, b, c) shared(d, e, u, v, w)
+ #pragma omp distribute lastprivate(d, e)
+ for (d = a; d < b; d++)
+ {
+ u[d] = v[d] + w[d];
+ e = c + d * 5;
+ }
+ return d + e;
+}
+
+__attribute__((noinline, noclone)) long
+f3 (long a1, long b1, long a2, long b2)
+{
+ long d1, d2;
+ #pragma omp target map(from: d1, d2)
+ #pragma omp teams default(none) shared(a1, b1, a2, b2, d1, d2, u, v, w)
+ #pragma omp distribute firstprivate (a1, b1, a2, b2) lastprivate(d1, d2) collapse(2)
+ for (d1 = a1; d1 < b1; d1++)
+ for (d2 = a2; d2 < b2; d2++)
+ u[d1 * 32 + d2] = v[d1 * 32 + d2] + w[d1 * 32 + d2];
+ return d1 + d2;
+}
+
+int
+main ()
+{
+ if (f2 (0, 1024, 17) != 1024 + (17 + 5 * 1023)
+ || f3 (0, 32, 0, 32) != 64)
+ __builtin_abort ();
+ return 0;
+}
--- libgomp/testsuite/libgomp.c++/pr66199-3.C.jj 2015-10-22 17:14:28.277980917 +0200
+++ libgomp/testsuite/libgomp.c++/pr66199-3.C 2015-10-22 17:14:28.276980932 +0200
@@ -0,0 +1,4 @@
+// PR middle-end/66199
+// { dg-do run }
+
+#include "../libgomp.c/pr66199-3.c"
--- libgomp/testsuite/libgomp.c++/pr66199-4.C.jj 2015-10-22 17:14:28.280980875 +0200
+++ libgomp/testsuite/libgomp.c++/pr66199-4.C 2015-10-22 17:14:28.279980889 +0200
@@ -0,0 +1,4 @@
+// PR middle-end/66199
+// { dg-do run }
+
+#include "../libgomp.c/pr66199-4.c"
--- libgomp/testsuite/libgomp.c++/pr66199-5.C.jj 2015-10-22 17:14:28.283980832 +0200
+++ libgomp/testsuite/libgomp.c++/pr66199-5.C 2015-10-22 17:14:28.282980846 +0200
@@ -0,0 +1,4 @@
+// PR middle-end/66199
+// { dg-do run }
+
+#include "../libgomp.c/pr66199-5.c"
--- libgomp/testsuite/libgomp.c++/pr66199-6.C.jj 2015-10-22 17:14:28.287980775 +0200
+++ libgomp/testsuite/libgomp.c++/pr66199-6.C 2015-10-22 17:14:28.285980803 +0200
@@ -0,0 +1,4 @@
+// PR middle-end/66199
+// { dg-do run }
+
+#include "../libgomp.c/pr66199-6.c"
--- libgomp/testsuite/libgomp.c++/pr66199-7.C.jj 2015-10-22 17:14:28.290980732 +0200
+++ libgomp/testsuite/libgomp.c++/pr66199-7.C 2015-10-22 17:14:28.289980746 +0200
@@ -0,0 +1,4 @@
+// PR middle-end/66199
+// { dg-do run }
+
+#include "../libgomp.c/pr66199-7.c"
--- libgomp/testsuite/libgomp.c++/pr66199-8.C.jj 2015-10-22 17:14:28.293980689 +0200
+++ libgomp/testsuite/libgomp.c++/pr66199-8.C 2015-10-22 17:14:28.292980703 +0200
@@ -0,0 +1,4 @@
+// PR middle-end/66199
+// { dg-do run }
+
+#include "../libgomp.c/pr66199-8.c"
--- libgomp/testsuite/libgomp.c++/pr66199-9.C.jj 2015-10-22 17:14:28.296980646 +0200
+++ libgomp/testsuite/libgomp.c++/pr66199-9.C 2015-10-22 17:14:28.295980661 +0200
@@ -0,0 +1,4 @@
+// PR middle-end/66199
+// { dg-do run }
+
+#include "../libgomp.c/pr66199-9.c"
Jakub
More information about the Gcc-patches
mailing list