This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
[gomp4] error on acc loops not associated with offloaded acc regions
- From: Cesar Philippidis <cesar at codesourcery dot com>
- To: "gcc-patches at gcc dot gnu dot org" <gcc-patches at gcc dot gnu dot org>
- Date: Mon, 28 Sep 2015 10:08:34 -0700
- Subject: [gomp4] error on acc loops not associated with offloaded acc regions
- Authentication-results: sourceware.org; auth=none
I've applied this patch to gomp-4_0-branch which teaches omplower how to
error when it detects acc loops which aren't nested inside an acc
parallel or kernels region or located within a function marked as an acc
routine. A couple of test cases needed to be updated.
The error message is kind of long. Let me know if it should be revised.
Cesar
2015-09-28 Cesar Philippidis <cesar@codesourcery.com>
gcc/
* omp-low.c (check_omp_nesting_restrictions): Check for acc loops not
associated with acc regions or routines.
gcc/testsuite/
* c-c++-common/goacc/non-routine.c: New test.
* c-c++-common/goacc-gomp/nesting-1.c: Add checks for invalid loop
nesting.
* c-c++-common/goacc-gomp/nesting-fail-1.c: Likewise.
* c-c++-common/goacc/clauses-fail.c: Likewise.
* c-c++-common/goacc/sb-1.c: Likewise.
* c-c++-common/goacc/sb-3.c: Likewise.
* gcc.dg/goacc/sb-1.c: Likewise.
* gcc.dg/goacc/sb-3.c: Likewise.
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 99b3939..2329a71 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -2901,6 +2901,14 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx)
}
return true;
}
+ if (is_gimple_omp_oacc (stmt) && ctx == NULL
+ && get_oacc_fn_attrib (current_function_decl) == NULL)
+ {
+ error_at (gimple_location (stmt),
+ "acc loops must be associated with an acc region or "
+ "routine");
+ return false;
+ }
/* FALLTHRU */
case GIMPLE_CALL:
if (is_gimple_call (stmt)
diff --git a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c
index b38e181..75d6a1d 100644
--- a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c
+++ b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c
@@ -20,6 +20,7 @@ f_acc_kernels (void)
}
}
+#pragma acc routine
void
f_acc_loop (void)
{
diff --git a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
index 14c6aa6..6d91484 100644
--- a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
+++ b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
@@ -361,72 +361,72 @@ f_acc_data (void)
void
f_acc_loop (void)
{
-#pragma acc loop
+#pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
for (i = 0; i < 2; ++i)
{
-#pragma omp parallel /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+#pragma omp parallel
;
}
-#pragma acc loop
+#pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
for (i = 0; i < 2; ++i)
{
-#pragma omp for /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+#pragma omp for
for (i = 0; i < 3; i++)
;
}
-#pragma acc loop
+#pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
for (i = 0; i < 2; ++i)
{
-#pragma omp sections /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+#pragma omp sections
{
;
}
}
-#pragma acc loop
+#pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
for (i = 0; i < 2; ++i)
{
-#pragma omp single /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+#pragma omp single
;
}
-#pragma acc loop
+#pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
for (i = 0; i < 2; ++i)
{
-#pragma omp task /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+#pragma omp task
;
}
-#pragma acc loop
+#pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
for (i = 0; i < 2; ++i)
{
-#pragma omp master /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+#pragma omp master
;
}
-#pragma acc loop
+#pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
for (i = 0; i < 2; ++i)
{
-#pragma omp critical /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+#pragma omp critical
;
}
-#pragma acc loop
+#pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
for (i = 0; i < 2; ++i)
{
-#pragma omp ordered /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+#pragma omp ordered
;
}
-#pragma acc loop
+#pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
for (i = 0; i < 2; ++i)
{
-#pragma omp target /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+#pragma omp target
;
-#pragma omp target data /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+#pragma omp target data
;
-#pragma omp target update to(i) /* { dg-error "non-OpenACC construct inside of OpenACC region" } */
+#pragma omp target update to(i)
}
}
diff --git a/gcc/testsuite/c-c++-common/goacc/clauses-fail.c b/gcc/testsuite/c-c++-common/goacc/clauses-fail.c
index 8990180..5a572f6 100644
--- a/gcc/testsuite/c-c++-common/goacc/clauses-fail.c
+++ b/gcc/testsuite/c-c++-common/goacc/clauses-fail.c
@@ -16,3 +16,5 @@ f (void)
for (i = 0; i < 2; ++i)
;
}
+
+/* { dg-error "acc loops must be associated with an acc region or routine" "" { target *-*-* } 15 } */
\ No newline at end of file
diff --git a/gcc/testsuite/c-c++-common/goacc/loop-1.c b/gcc/testsuite/c-c++-common/goacc/loop-1.c
index 5e1a248..bb8b9f3 100644
--- a/gcc/testsuite/c-c++-common/goacc/loop-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/loop-1.c
@@ -36,16 +36,16 @@ int test1()
i = d;
a[i] = 1;
}
- #pragma acc loop
+ #pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
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 "acc loops must be associated with an acc region or routine" } */
for (i = 1; i < 10; i++)
{
}
- #pragma acc loop
+ #pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
for (i = 1; i < 10; i+=2)
{
a[i] = i;
diff --git a/gcc/testsuite/c-c++-common/goacc/non-routine.c b/gcc/testsuite/c-c++-common/goacc/non-routine.c
new file mode 100644
index 0000000..0af69cf
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/non-routine.c
@@ -0,0 +1,16 @@
+/* This program validates the behavior of acc loops which are
+ not associated with a parallel or kernles region or routine. */
+
+/* { dg-do compile } */
+
+int
+main ()
+{
+ int i, v = 0;
+
+#pragma acc loop gang reduction (+:v) /* { dg-error "acc loops must be associated with an acc region or routine" } */
+ for (i = 0; i < 10; i++)
+ v++;
+
+ return v;
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/sb-1.c b/gcc/testsuite/c-c++-common/goacc/sb-1.c
index 5e55c95..968ce5f 100644
--- a/gcc/testsuite/c-c++-common/goacc/sb-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/sb-1.c
@@ -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 "acc loops must be associated with an acc region or routine" } */
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 "acc loops must be associated with an acc region or routine" } */
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 "acc loops must be associated with an acc region or routine" } */
for (l = 0; l < 2; ++l)
{
int i;
diff --git a/gcc/testsuite/c-c++-common/goacc/sb-3.c b/gcc/testsuite/c-c++-common/goacc/sb-3.c
index 1567a10..dd57a8f 100644
--- a/gcc/testsuite/c-c++-common/goacc/sb-3.c
+++ b/gcc/testsuite/c-c++-common/goacc/sb-3.c
@@ -3,7 +3,7 @@
void f (void)
{
int i, j;
-#pragma acc loop
+#pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
for(i = 1; i < 30; i++)
{
if (i == 7) goto out; // { dg-error "invalid branch to/from OpenACC structured block" }
diff --git a/gcc/testsuite/gcc.dg/goacc/sb-1.c b/gcc/testsuite/gcc.dg/goacc/sb-1.c
index 0a3316c..cf61518 100644
--- a/gcc/testsuite/gcc.dg/goacc/sb-1.c
+++ b/gcc/testsuite/gcc.dg/goacc/sb-1.c
@@ -9,7 +9,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 "acc loops must be associated with an acc region or routine" } */
for (l = 0; l < 2; ++l)
goto bad1; // { dg-error "invalid branch to/from OpenACC structured block" }
@@ -32,7 +32,7 @@ void foo()
}
goto bad2_loop; // { dg-error "invalid entry to OpenACC structured block" }
- #pragma acc loop
+ #pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
for (l = 0; l < 2; ++l)
{
bad2_loop: ;
@@ -62,7 +62,7 @@ void foo()
{ ok1_data: break; }
}
- #pragma acc loop
+ #pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
for (l = 0; l < 2; ++l)
{
int i;
diff --git a/gcc/testsuite/gcc.dg/goacc/sb-3.c b/gcc/testsuite/gcc.dg/goacc/sb-3.c
index acbe9d5..7999ebe 100644
--- a/gcc/testsuite/gcc.dg/goacc/sb-3.c
+++ b/gcc/testsuite/gcc.dg/goacc/sb-3.c
@@ -1,7 +1,7 @@
void f (void)
{
int i, j;
-#pragma acc loop
+#pragma acc loop /* { dg-error "acc loops must be associated with an acc region or routine" } */
for(i = 1; i < 30; i++)
{
if (i == 7) goto out; // { dg-error "invalid branch to/from OpenACC structured block" }