This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
[OpenACC] loop nesting check
- From: Nathan Sidwell <nathan at acm dot org>
- To: GCC Patches <gcc-patches at gcc dot gnu dot org>
- Cc: Jakub Jelinek <jakub at redhat dot com>, Bernd Schmidt <bschmidt at redhat dot com>
- Date: Wed, 21 Oct 2015 11:59:03 -0400
- Subject: [OpenACC] loop nesting check
- Authentication-results: sourceware.org; auth=none
Hi,
in preparing a patch set of the OpenACC execution model for trunk, I discovered
some broken tests in the testsuite. We were failing to diagnose some
incorrectly nested loops, and that led to ICEs in my patch set.
This patch implements a check for openacc loop directives, checking that they
are immediately inside one of:
1) another openacc loop
2) an openacc offload region
3) an openacc routine
The broken tests are amended with the now expected diagnostic.
tested on x86_64-linux-gnu with nvptx accelerator.
ok for trunk?
nathan
2015-10-21 Nathan Sidwell <nathan@codesourcery.com>
gcc/
* omp-low.c (check_omp_nesting_restrictions): Check OpenACC loop
nesting.
testsuite/
* c-c++-common/goacc/clauses-fail.c: Adjust errors.
* c-c++-common/goacc/sb-1.c: Adjust errors.
* c-c++-common/goacc/sb-3.c: Adjust errors.
* c-c++-common/goacc/loop-1.c: Adjust errors.
* c-c++-common/goacc/nesting-1.c: Adjust errors.
* c-c++-common/goacc-gomp/nesting-fail-1.c: Adjust errors.
* c-c++-common/goacc-gomp/nesting-1.c: Adjust errors.
Index: gcc/omp-low.c
===================================================================
--- gcc/omp-low.c (revision 229101)
+++ gcc/omp-low.c (working copy)
@@ -3178,6 +3200,43 @@ check_omp_nesting_restrictions (gimple *
/* We split taskloop into task and nested taskloop in it. */
if (gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_TASKLOOP)
return true;
+ if (gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_OACC_LOOP)
+ {
+ bool ok = false;
+
+ if (ctx)
+ switch (gimple_code (ctx->stmt))
+ {
+ case GIMPLE_OMP_FOR:
+ ok = (gimple_omp_for_kind (ctx->stmt)
+ == GF_OMP_FOR_KIND_OACC_LOOP);
+ break;
+
+ case GIMPLE_OMP_TARGET:
+ switch (gimple_omp_target_kind (ctx->stmt))
+ {
+ case GF_OMP_TARGET_KIND_OACC_PARALLEL:
+ case GF_OMP_TARGET_KIND_OACC_KERNELS:
+ ok = true;
+ break;
+
+ default:
+ break;
+ }
+
+ default:
+ break;
+ }
+ else if (get_oacc_fn_attrib (current_function_decl))
+ ok = true;
+ if (!ok)
+ {
+ error_at (gimple_location (stmt),
+ "OpenACC loop directive must be associated with"
+ " an OpenACC compute region");
+ return false;
+ }
+ }
/* FALLTHRU */
case GIMPLE_CALL:
if (is_gimple_call (stmt)
Index: gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c (revision 229101)
+++ gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c (working copy)
@@ -5,7 +5,7 @@ f_omp_parallel (void)
{
int i;
-#pragma acc loop
+#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
for (i = 0; i < 2; ++i)
;
}
Index: gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c (revision 229101)
+++ gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c (working copy)
@@ -28,7 +28,7 @@ f_omp (void)
#pragma acc update host(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
#pragma acc enter data copyin(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
#pragma acc exit data delete(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
-#pragma acc loop /* { dg-error "may not be closely nested" } */
+#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
for (i = 0; i < 2; ++i)
;
}
@@ -63,7 +63,7 @@ f_omp (void)
}
#pragma omp section
{
-#pragma acc loop /* { dg-error "may not be closely nested" } */
+#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
for (i = 0; i < 2; ++i)
;
}
@@ -80,7 +80,7 @@ f_omp (void)
#pragma acc update host(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
#pragma acc enter data copyin(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
#pragma acc exit data delete(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
-#pragma acc loop /* { dg-error "may not be closely nested" } */
+#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
for (i = 0; i < 2; ++i)
;
}
@@ -96,7 +96,7 @@ f_omp (void)
#pragma acc update host(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
#pragma acc enter data copyin(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
#pragma acc exit data delete(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
-#pragma acc loop /* { dg-error "may not be closely nested" } */
+#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
for (i = 0; i < 2; ++i)
;
}
@@ -112,7 +112,7 @@ f_omp (void)
#pragma acc update host(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
#pragma acc enter data copyin(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
#pragma acc exit data delete(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
-#pragma acc loop /* { dg-error "may not be closely nested" } */
+#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
for (i = 0; i < 2; ++i)
;
}
@@ -128,7 +128,7 @@ f_omp (void)
#pragma acc update host(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
#pragma acc enter data copyin(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
#pragma acc exit data delete(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
-#pragma acc loop /* { dg-error "may not be closely nested" } */
+#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
for (i = 0; i < 2; ++i)
;
}
@@ -144,7 +144,7 @@ f_omp (void)
#pragma acc update host(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
#pragma acc enter data copyin(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
#pragma acc exit data delete(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */
-#pragma acc loop /* { dg-error "may not be closely nested" } */
+#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
for (i = 0; i < 2; ++i)
;
}
@@ -160,7 +160,7 @@ f_omp (void)
#pragma acc update host(i) /* { dg-error "OpenACC update construct inside of OpenMP target region" } */
#pragma acc enter data copyin(i) /* { dg-error "OpenACC enter/exit data construct inside of OpenMP target region" } */
#pragma acc exit data delete(i) /* { dg-error "OpenACC enter/exit data construct inside of OpenMP target region" } */
-#pragma acc loop
+#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
for (i = 0; i < 2; ++i)
;
}
@@ -379,6 +379,7 @@ f_acc_data (void)
void
f_acc_loop (void)
{
+#pragma acc parallel
#pragma acc loop
for (i = 0; i < 2; ++i)
{
@@ -386,6 +387,7 @@ f_acc_loop (void)
;
}
+#pragma acc parallel
#pragma acc loop
for (i = 0; i < 2; ++i)
{
@@ -394,6 +396,7 @@ f_acc_loop (void)
;
}
+#pragma acc parallel
#pragma acc loop
for (i = 0; i < 2; ++i)
{
@@ -403,6 +406,7 @@ f_acc_loop (void)
}
}
+#pragma acc parallel
#pragma acc loop
for (i = 0; i < 2; ++i)
{
@@ -410,6 +414,7 @@ f_acc_loop (void)
;
}
+#pragma acc parallel
#pragma acc loop
for (i = 0; i < 2; ++i)
{
@@ -417,6 +422,7 @@ f_acc_loop (void)
;
}
+#pragma acc parallel
#pragma acc loop
for (i = 0; i < 2; ++i)
{
@@ -424,6 +430,7 @@ f_acc_loop (void)
;
}
+#pragma acc parallel
#pragma acc loop
for (i = 0; i < 2; ++i)
{
@@ -431,6 +438,7 @@ f_acc_loop (void)
;
}
+#pragma acc parallel
#pragma acc loop
for (i = 0; i < 2; ++i)
{
@@ -438,6 +446,7 @@ f_acc_loop (void)
i = 0; /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
}
+#pragma acc parallel
#pragma acc loop
for (i = 0; i < 2; ++i)
{
@@ -445,6 +454,7 @@ f_acc_loop (void)
;
}
+#pragma acc parallel
#pragma acc loop
for (i = 0; i < 2; ++i)
{
Index: gcc/testsuite/c-c++-common/goacc/nesting-1.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/nesting-1.c (revision 229101)
+++ gcc/testsuite/c-c++-common/goacc/nesting-1.c (working copy)
@@ -58,7 +58,7 @@ f_acc_data (void)
#pragma acc exit data delete(i)
-#pragma acc loop
+#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
for (i = 0; i < 2; ++i)
;
@@ -93,7 +93,7 @@ f_acc_data (void)
#pragma acc exit data delete(i)
-#pragma acc loop
+#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
for (i = 0; i < 2; ++i)
;
}
Index: gcc/testsuite/c-c++-common/goacc/loop-1.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/loop-1.c (revision 229101)
+++ gcc/testsuite/c-c++-common/goacc/loop-1.c (working copy)
@@ -38,16 +38,16 @@ int test1()
i = d;
a[i] = 1;
}
- #pragma acc loop
+ #pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
for (i = 1; i < 30; i++ )
if (i == 16) break; /* { dg-error "break statement used" } */
/* different types of for loop are allowed */
- #pragma acc loop
+ #pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
for (i = 1; i < 10; i++)
{
}
- #pragma acc loop
+ #pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
for (i = 1; i < 10; i+=2)
{
a[i] = i;
Index: gcc/testsuite/c-c++-common/goacc/clauses-fail.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/clauses-fail.c (revision 229101)
+++ gcc/testsuite/c-c++-common/goacc/clauses-fail.c (working copy)
@@ -12,6 +12,7 @@ f (void)
#pragma acc data two /* { dg-error "expected '#pragma acc' clause before 'two'" } */
;
+#pragma acc parallel
#pragma acc loop deux /* { dg-error "expected '#pragma acc' clause before 'deux'" } */
for (i = 0; i < 2; ++i)
;
Index: gcc/testsuite/c-c++-common/goacc/sb-1.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/sb-1.c (revision 229101)
+++ gcc/testsuite/c-c++-common/goacc/sb-1.c (working copy)
@@ -11,7 +11,7 @@ void foo()
goto bad1; // { dg-error "invalid branch to/from OpenACC structured block" }
#pragma acc data
goto bad1; // { dg-error "invalid branch to/from OpenACC structured block" }
- #pragma acc loop
+ #pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
for (l = 0; l < 2; ++l)
goto bad1; // { dg-error "invalid branch to/from OpenACC structured block" }
@@ -34,7 +34,7 @@ void foo()
}
goto bad2_loop; // { dg-error "invalid entry to OpenACC structured block" }
- #pragma acc loop
+ #pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
for (l = 0; l < 2; ++l)
{
bad2_loop: ;
@@ -64,7 +64,7 @@ void foo()
{ ok1_data: break; }
}
- #pragma acc loop
+ #pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
for (l = 0; l < 2; ++l)
{
int i;
Index: gcc/testsuite/c-c++-common/goacc/sb-3.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/sb-3.c (revision 229101)
+++ gcc/testsuite/c-c++-common/goacc/sb-3.c (working copy)
@@ -3,11 +3,11 @@
void f (void)
{
int i, j;
-#pragma acc loop
+#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
for(i = 1; i < 30; i++)
{
if (i == 7) goto out; // { dg-error "invalid branch to/from OpenACC structured block" }
-#pragma acc loop // { dg-error "work-sharing region may not be closely nested inside of work-sharing, critical, ordered, master or explicit task region" }
+#pragma acc loop
for(j = 5; j < 10; j++)
{
if (i == 6 && j == 7) goto out; // { dg-error "invalid branch to/from OpenACC structured block" }