Hi! Currently if a target* pragma appears within a target region, GCC successfully compiles such code (with a warning). But the binary fails at run-time, since it tries to call GOMP_target* functions on target.
The spec says: "If a target, target update, or target data construct appears within a target region then the behavior is unspecified." I see 2 options to make the behavior more user-friendly: 1. To return an error at compile-time. 2. To check at run-time in libgomp whether GOMP_target* is called on target, and perform target-fallback if so. If we will select option #1, the patch is ready. diff --git a/gcc/omp-low.c b/gcc/omp-low.c index eaad52a..ae8b90a 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -2682,8 +2682,14 @@ check_omp_nesting_restrictions (gimple stmt, omp_context *ctx) } break; case GIMPLE_OMP_TARGET: - for (; ctx != NULL; ctx = ctx->outer) - if (is_targetreg_ctx (ctx)) + { + bool target_p = false; + if (cgraph_node::get (current_function_decl)->offloadable) + target_p = true; + for (; ctx != NULL && !target_p; ctx = ctx->outer) + if (is_targetreg_ctx (ctx)) + target_p = true; + if (target_p) { const char *name; switch (gimple_omp_target_kind (stmt)) @@ -2693,9 +2699,11 @@ check_omp_nesting_restrictions (gimple stmt, omp_context *ctx) case GF_OMP_TARGET_KIND_UPDATE: name = "target update"; break; default: gcc_unreachable (); } - warning_at (gimple_location (stmt), 0, - "%s construct inside of target region", name); + error_at (gimple_location (stmt), + "%s construct inside of target region", name); + return false; } + } break; default: break; diff --git a/gcc/testsuite/c-c++-common/gomp/nesting-error-1.c b/gcc/testsuite/c-c++-common/gomp/nesting-error-1.c new file mode 100644 index 0000000..ff6a75e --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/nesting-error-1.c @@ -0,0 +1,23 @@ +extern int i; + +void +f_omp_target (void) +{ +#pragma omp target + { +#pragma omp target /* { dg-error "target construct inside of target region" } */ + ; +#pragma omp target data /* { dg-error "target data construct inside of target region" } */ + ; +#pragma omp target update to(i) /* { dg-error "target update construct inside of target region" } */ + +#pragma omp parallel + { +#pragma omp target /* { dg-error "target construct inside of target region" } */ + ; +#pragma omp target data /* { dg-error "target data construct inside of target region" } */ + ; +#pragma omp target update to(i) /* { dg-error "target update construct inside of target region" } */ + } + } +} diff --git a/gcc/testsuite/c-c++-common/gomp/nesting-warn-1.c b/gcc/testsuite/c-c++-common/gomp/nesting-warn-1.c deleted file mode 100644 index c39dd49..0000000 --- a/gcc/testsuite/c-c++-common/gomp/nesting-warn-1.c +++ /dev/null @@ -1,23 +0,0 @@ -extern int i; - -void -f_omp_target (void) -{ -#pragma omp target - { -#pragma omp target /* { dg-warning "target construct inside of target region" } */ - ; -#pragma omp target data /* { dg-warning "target data construct inside of target region" } */ - ; -#pragma omp target update to(i) /* { dg-warning "target update construct inside of target region" } */ - -#pragma omp parallel - { -#pragma omp target /* { dg-warning "target construct inside of target region" } */ - ; -#pragma omp target data /* { dg-warning "target data construct inside of target region" } */ - ; -#pragma omp target update to(i) /* { dg-warning "target update construct inside of target region" } */ - } - } -} diff --git a/gcc/testsuite/gfortran.dg/gomp/target3.f90 b/gcc/testsuite/gfortran.dg/gomp/target3.f90 index 53a9682..e451548 100644 --- a/gcc/testsuite/gfortran.dg/gomp/target3.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/target3.f90 @@ -4,9 +4,18 @@ subroutine foo (r) integer :: i, r !$omp target - !$omp target teams distribute parallel do reduction (+: r) ! { dg-warning "target construct inside of target region" } + !$omp target teams distribute parallel do reduction (+: r) ! { dg-error "target construct inside of target region" } do i = 1, 10 r = r + 1 end do !$omp end target end subroutine + +subroutine bar (r) + !$omp declare target + integer :: i, r + !$omp target teams distribute parallel do reduction (+: r) ! { dg-error "target construct inside of target region" } + do i = 1, 10 + r = r + 1 + end do +end subroutine -- Ilya