diff options
Diffstat (limited to 'gcc')
| -rw-r--r-- | gcc/ChangeLog | 5 | ||||
| -rw-r--r-- | gcc/omp-low.c | 41 | ||||
| -rw-r--r-- | gcc/testsuite/ChangeLog | 6 | ||||
| -rw-r--r-- | gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c | 8 | ||||
| -rw-r--r-- | gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c | 8 |
5 files changed, 55 insertions, 13 deletions
diff --git a/gcc/ChangeLog b/gcc/ChangeLog index ea4e2d5..5c23836 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,8 @@ +2016-03-05 Tom de Vries <tom@codesourcery.com> + + * omp-low.c (check_omp_nesting_restrictions): Check for non-oacc + construct in oacc routine. Check for oacc region in oacc routine. + 2016-03-04 Jakub Jelinek <jakub@redhat.com> PR target/70062 diff --git a/gcc/omp-low.c b/gcc/omp-low.c index ecbf74a..c69fe44 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -3236,19 +3236,26 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx) /* No nesting of non-OpenACC STMT (that is, an OpenMP one, or a GOMP builtin) inside an OpenACC CTX. */ if (!(is_gimple_omp (stmt) - && is_gimple_omp_oacc (stmt))) - { - for (omp_context *octx = ctx; octx != NULL; octx = octx->outer) - if (is_gimple_omp (octx->stmt) - && is_gimple_omp_oacc (octx->stmt) - /* Except for atomic codes that we share with OpenMP. */ - && ! (gimple_code (stmt) == GIMPLE_OMP_ATOMIC_LOAD - || gimple_code (stmt) == GIMPLE_OMP_ATOMIC_STORE)) - { - error_at (gimple_location (stmt), - "non-OpenACC construct inside of OpenACC region"); - return false; - } + && is_gimple_omp_oacc (stmt)) + /* Except for atomic codes that we share with OpenMP. */ + && !(gimple_code (stmt) == GIMPLE_OMP_ATOMIC_LOAD + || gimple_code (stmt) == GIMPLE_OMP_ATOMIC_STORE)) + { + if (get_oacc_fn_attrib (cfun->decl) != NULL) + { + error_at (gimple_location (stmt), + "non-OpenACC construct inside of OpenACC routine"); + return false; + } + else + for (omp_context *octx = ctx; octx != NULL; octx = octx->outer) + if (is_gimple_omp (octx->stmt) + && is_gimple_omp_oacc (octx->stmt)) + { + error_at (gimple_location (stmt), + "non-OpenACC construct inside of OpenACC region"); + return false; + } } if (ctx != NULL) @@ -3715,6 +3722,14 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx) kind == OMP_CLAUSE_DEPEND_SOURCE ? "source" : "sink"); return false; } + if (is_gimple_omp_offloaded (stmt) + && get_oacc_fn_attrib (cfun->decl) != NULL) + { + error_at (gimple_location (stmt), + "OpenACC region inside of OpenACC routine, nested " + "parallelism not supported yet"); + return false; + } for (; ctx != NULL; ctx = ctx->outer) { if (gimple_code (ctx->stmt) != GIMPLE_OMP_TARGET) diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog index ca65abf..51469ea 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,3 +1,9 @@ +2016-03-05 Tom de Vries <tom@codesourcery.com> + + * c-c++-common/goacc/nesting-fail-1.c (f_acc_routine): New function. + * c-c++-common/goacc-gomp/nesting-fail-1.c (f_acc_routine): New + function. + 2016-03-05 Patrick Palka <ppalka@gcc.gnu.org> PR c++/66786 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 1a44721..5e3f183 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 @@ -439,3 +439,11 @@ f_acc_loop (void) #pragma omp target update to(i) /* { dg-error "non-OpenACC construct inside of OpenACC region" } */ } } + +#pragma acc routine +void +f_acc_routine (void) +{ +#pragma omp target /* { dg-error "non-OpenACC construct inside of OpenACC routine" } */ + ; +} diff --git a/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c b/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c index 7a36074..506a1ae 100644 --- a/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c +++ b/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c @@ -37,3 +37,11 @@ f_acc_kernels (void) #pragma acc exit data delete(i) /* { dg-error ".enter/exit data. construct inside of .kernels. region" } */ } } + +#pragma acc routine +void +f_acc_routine (void) +{ +#pragma acc parallel /* { dg-error "OpenACC region inside of OpenACC routine, nested parallelism not supported yet" } */ + ; +} |