From: Nathan Sidwell Date: Wed, 21 Oct 2015 16:14:01 +0000 (+0000) Subject: omp-low.c (check_omp_nesting_restrictions): Check OpenACC loop nesting. X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=68d58afb65a42cb7886552a1532fc11dc06eaed4;p=gcc.git omp-low.c (check_omp_nesting_restrictions): Check OpenACC loop nesting. 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. From-SVN: r229129 --- diff --git a/gcc/ChangeLog b/gcc/ChangeLog index d5296ae5db4..21164b13ca3 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,8 @@ +2015-10-21 Nathan Sidwell + + * omp-low.c (check_omp_nesting_restrictions): Check OpenACC loop + nesting. + 2015-10-21 Ilya Enkovich * doc/tm.texi: Regenerated. diff --git a/gcc/omp-low.c b/gcc/omp-low.c index b71609b76a4..ad7c017affb 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -3178,6 +3178,43 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx) /* 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) diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog index 3f325b9ad3c..db0f89665a2 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,3 +1,13 @@ +2015-10-21 Nathan Sidwell + + * 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. + 2015-10-21 Ilya Enkovich * g++.dg/ext/vector22.C: Allow VEC_COND_EXPR. 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 df45bcf05fc..1c17818e415 100644 --- a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c +++ b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c @@ -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) ; } 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 ac614e7d646..0c8ea54d53d 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 @@ -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) { diff --git a/gcc/testsuite/c-c++-common/goacc/clauses-fail.c b/gcc/testsuite/c-c++-common/goacc/clauses-fail.c index 899018026d3..661d364f066 100644 --- a/gcc/testsuite/c-c++-common/goacc/clauses-fail.c +++ b/gcc/testsuite/c-c++-common/goacc/clauses-fail.c @@ -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) ; diff --git a/gcc/testsuite/c-c++-common/goacc/loop-1.c b/gcc/testsuite/c-c++-common/goacc/loop-1.c index fea40e0ab61..2d8f4d5a8d3 100644 --- a/gcc/testsuite/c-c++-common/goacc/loop-1.c +++ b/gcc/testsuite/c-c++-common/goacc/loop-1.c @@ -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; diff --git a/gcc/testsuite/c-c++-common/goacc/nesting-1.c b/gcc/testsuite/c-c++-common/goacc/nesting-1.c index b4b863fb860..3a8f838906c 100644 --- a/gcc/testsuite/c-c++-common/goacc/nesting-1.c +++ b/gcc/testsuite/c-c++-common/goacc/nesting-1.c @@ -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) ; } diff --git a/gcc/testsuite/c-c++-common/goacc/sb-1.c b/gcc/testsuite/c-c++-common/goacc/sb-1.c index 5e55c9516f2..1ce41df0c09 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 "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; diff --git a/gcc/testsuite/c-c++-common/goacc/sb-3.c b/gcc/testsuite/c-c++-common/goacc/sb-3.c index 147b7b0e845..620498ecd62 100644 --- a/gcc/testsuite/c-c++-common/goacc/sb-3.c +++ b/gcc/testsuite/c-c++-common/goacc/sb-3.c @@ -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" }