@@ -1,5 +1,8 @@
2014-11-12 Thomas Schwinge <thomas@codesourcery.com>
+ * gimplify.c (gimplify_omp_for): Eliminate govd_private; always
+ use GOVD_PRIVATE.
+
* omp-low.c (scan_oacc_offload, expand_oacc_offload)
(lower_oacc_offload): Merge into scan_omp_target,
expand_omp_target, lower_omp_target, respectively. Update all
@@ -6843,7 +6843,6 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
gimple_seq for_body, for_pre_body;
int i;
bool simd;
- enum gimplify_omp_var_data govd_private;
enum omp_region_type ort;
bitmap has_decl_expr = NULL;
@@ -6855,18 +6854,15 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
case CILK_FOR:
case OMP_DISTRIBUTE:
simd = false;
- govd_private = GOVD_PRIVATE;
ort = ORT_WORKSHARE;
break;
case OMP_SIMD:
case CILK_SIMD:
simd = true;
- govd_private = GOVD_PRIVATE;
ort = ORT_SIMD;
break;
case OACC_LOOP:
simd = false;
- govd_private = /* TODO */ GOVD_LOCAL;
ort = /* TODO */ ORT_WORKSHARE;
break;
default:
@@ -6928,7 +6924,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
gcc_assert (INTEGRAL_TYPE_P (TREE_TYPE (decl))
|| POINTER_TYPE_P (TREE_TYPE (decl)));
- /* Make sure the iteration variable is some kind of private. */
+ /* Make sure the iteration variable is private. */
tree c = NULL_TREE;
tree c2 = NULL_TREE;
if (orig_for_stmt != for_stmt)
@@ -6957,7 +6953,6 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
}
else
{
- gcc_assert (govd_private == GOVD_PRIVATE);
bool lastprivate
= (!has_decl_expr
|| !bitmap_bit_p (has_decl_expr, DECL_UID (decl)));
@@ -7001,7 +6996,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
else if (omp_is_private (gimplify_omp_ctxp, decl, 0))
omp_notice_variable (gimplify_omp_ctxp, decl, true);
else
- omp_add_variable (gimplify_omp_ctxp, decl, govd_private | GOVD_SEEN);
+ omp_add_variable (gimplify_omp_ctxp, decl, GOVD_PRIVATE | GOVD_SEEN);
/* If DECL is not a gimple register, create a temporary variable to act
as an iteration counter. This is valid, since DECL cannot be
@@ -7036,7 +7031,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
}
else
omp_add_variable (gimplify_omp_ctxp, var,
- govd_private | GOVD_SEEN);
+ GOVD_PRIVATE | GOVD_SEEN);
}
else
var = decl;
@@ -7193,7 +7188,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
t = TREE_VEC_ELT (OMP_FOR_INIT (for_stmt), i);
decl = TREE_OPERAND (t, 0);
var = create_tmp_var (TREE_TYPE (decl), get_name (decl));
- omp_add_variable (gimplify_omp_ctxp, var, govd_private | GOVD_SEEN);
+ omp_add_variable (gimplify_omp_ctxp, var, GOVD_PRIVATE | GOVD_SEEN);
TREE_OPERAND (t, 0) = var;
t = TREE_VEC_ELT (OMP_FOR_INCR (for_stmt), i);
TREE_OPERAND (t, 1) = copy_node (TREE_OPERAND (t, 1));
@@ -1,3 +1,7 @@
+2014-11-12 Thomas Schwinge <thomas@codesourcery.com>
+
+ * c-c++-common/goacc/loop-private-1.c: New file.
+
2014-11-07 Thomas Schwinge <thomas@codesourcery.com>
James Norris <jnorris@codesourcery.com>
new file mode 100644
@@ -0,0 +1,14 @@
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+void
+f (int i, int j)
+{
+#pragma acc kernels
+#pragma acc loop collapse(2)
+ for (i = 0; i < 20; ++i)
+ for (j = 0; j < 25; ++j)
+ ;
+}
+
+/* { dg-final { scan-tree-dump-times "#pragma acc loop collapse\\(2\\) private\\(j\\) private\\(i\\)" 1 "gimple" } } */
+/* { dg-final { cleanup-tree-dump "gimple" } } */
@@ -1,3 +1,11 @@
+2014-11-12 Thomas Schwinge <thomas@codesourcery.com>
+
+ * testsuite/libgomp.oacc-c-c++-common/collapse-4.c: New file.
+ * testsuite/libgomp.oacc-c/collapse-4.c: Turn into an execution
+ test, check result.
+ * testsuite/libgomp.oacc-c-c++-common/data-2.c: Update.
+ * testsuite/libgomp.oacc-c-c++-common/data-3.c: Likewise.
+
2014-11-07 Thomas Schwinge <thomas@codesourcery.com>
* plugin-nvptx.c: Remove.
similarity index 56%
copy from libgomp/testsuite/libgomp.oacc-c/collapse-4.c
copy to libgomp/testsuite/libgomp.oacc-c-c++-common/collapse-4.c
@@ -1,6 +1,4 @@
-/* TODO: change to a run test once libgomp supports all data clauses. */
-/* { dg-do compile } */
-/* { dg-options "-O2 -std=c99" } */
+/* See also ../libgomp.oacc-c/collapse-4.c. */
#include <string.h>
@@ -9,17 +7,21 @@ main (void)
{
int l = 0;
int b[3][3];
+ int i, j;
memset (b, '\0', sizeof (b));
#pragma acc parallel copy(b[0:3][0:3]) copy(l)
{
#pragma acc loop collapse(2) reduction(+:l)
- for (int i = 0; i < 2; i++)
- for (int j = 0; j < 2; j++)
+ for (i = 0; i < 2; i++)
+ for (j = 0; j < 2; j++)
if (b[i][j] != 16)
l += 1;
}
+ if (l != 2 * 2)
+ __builtin_abort();
+
return 0;
}
@@ -28,8 +28,8 @@ main (int argc, char **argv)
#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async
#pragma acc parallel async wait
#pragma acc loop
- for (int ii = 0; ii < N; ii++)
- b[ii] = a[ii];
+ for (i = 0; i < N; i++)
+ b[i] = a[i];
#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) wait async
#pragma acc wait
@@ -52,8 +52,8 @@ main (int argc, char **argv)
#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async (1)
#pragma acc parallel async (1)
#pragma acc loop
- for (int ii = 0; ii < N; ii++)
- b[ii] = a[ii];
+ for (i = 0; i < N; i++)
+ b[i] = a[i];
#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) wait (1) async (1)
#pragma acc wait (1)
@@ -79,18 +79,18 @@ main (int argc, char **argv)
#pragma acc parallel async (1) wait (1)
#pragma acc loop
- for (int ii = 0; ii < N; ii++)
- b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
+ for (i = 0; i < N; i++)
+ b[i] = (a[i] * a[i] * a[i]) / a[i];
#pragma acc parallel async (2) wait (1)
#pragma acc loop
- for (int ii = 0; ii < N; ii++)
- c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
+ for (i = 0; i < N; i++)
+ c[i] = (a[i] + a[i] + a[i] + a[i]) / a[i];
#pragma acc parallel async (3) wait (1)
#pragma acc loop
- for (int ii = 0; ii < N; ii++)
- d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
+ for (i = 0; i < N; i++)
+ d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i];
#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) copyout (d[0:N]) wait (1, 2, 3) async (1)
#pragma acc wait (1)
@@ -28,8 +28,8 @@ main (int argc, char **argv)
#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async
#pragma acc parallel async wait
#pragma acc loop
- for (int ii = 0; ii < N; ii++)
- b[ii] = a[ii];
+ for (i = 0; i < N; i++)
+ b[i] = a[i];
#pragma acc update host (a[0:N], b[0:N]) async wait
#pragma acc wait
@@ -52,8 +52,8 @@ main (int argc, char **argv)
#pragma acc update device (a[0:N], b[0:N]) async (1)
#pragma acc parallel async (1)
#pragma acc loop
- for (int ii = 0; ii < N; ii++)
- b[ii] = a[ii];
+ for (i = 0; i < N; i++)
+ b[i] = a[i];
#pragma acc update host (a[0:N], b[0:N]) async (1) wait (1)
#pragma acc wait (1)
@@ -81,18 +81,18 @@ main (int argc, char **argv)
#pragma acc parallel async (1) wait (1,2)
#pragma acc loop
- for (int ii = 0; ii < N; ii++)
- b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
+ for (i = 0; i < N; i++)
+ b[i] = (a[i] * a[i] * a[i]) / a[i];
#pragma acc parallel async (2) wait (1,3)
#pragma acc loop
- for (int ii = 0; ii < N; ii++)
- c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
+ for (i = 0; i < N; i++)
+ c[i] = (a[i] + a[i] + a[i] + a[i]) / a[i];
#pragma acc parallel async (3) wait (1,3)
#pragma acc loop
- for (int ii = 0; ii < N; ii++)
- d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
+ for (i = 0; i < N; i++)
+ d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i];
#pragma acc update host (a[0:N], b[0:N], c[0:N], d[0:N]) async (1) wait (1,2,3)
#pragma acc wait (1)
@@ -1,5 +1,4 @@
-/* TODO: change to a run test once libgomp supports all data clauses. */
-/* { dg-do compile } */
+/* See also ../libgomp.oacc-c-c++-common/collapse-4.c. */
/* { dg-options "-O2 -std=c99" } */
#include <string.h>
@@ -21,5 +20,8 @@ main (void)
l += 1;
}
+ if (l != 2 * 2)
+ __builtin_abort();
+
return 0;
}