* [Patch] OpenMP/Fortran: Reject not strictly nested target -> teams [PR110725, PR71065] @ 2023-07-24 19:43 Tobias Burnus 2023-07-24 19:49 ` Jakub Jelinek 0 siblings, 1 reply; 6+ messages in thread From: Tobias Burnus @ 2023-07-24 19:43 UTC (permalink / raw) To: gcc-patches, fortran; +Cc: Jakub Jelinek [-- Attachment #1: Type: text/plain, Size: 909 bytes --] This patch adds diagnostic for additional code alongside a nested teams in a target region. The diagnostic is happening soon after parsing such that expressions in clauses are not yet expanded - those would end up before TEAMS and can be very complicated (e.g. assume an allocatable-returning function). (The patch diagnoses it in openmp.cc; after trans-openmp.cc it would already be to late.) Comments, remarks, suggestions? Tobias PS: Something similar is also needed for C/C++ but there templates and lambda functions might make it harder to implement. In any case, it has to be done in the FE. Tracked at PR71065 ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 [-- Attachment #2: omp-target-teams-diag.diff --] [-- Type: text/x-patch, Size: 12035 bytes --] OpenMP/Fortran: Reject not strictly nested target -> teams [PR110725, PR71065] OpenMP requires: "If a teams region is nested inside a target region, the corresponding target construct must not contain any statements, declarations or directives outside of the corresponding teams construct." Test for it! PR fortran/110725 PR middle-end/71065 gcc/fortran/ChangeLog: * gfortran.h (gfc_omp_clauses): Add contains_teams_construct. * openmp.cc (resolve_omp_target): New; check for teams nesting. (gfc_resolve_omp_directive): Call it. * parse.cc (decode_omp_directive): Set contains_teams_construct on enclosing ST_OMP_TARGET. gcc/testsuite/ChangeLog: * gfortran.dg/gomp/pr99226.f90: Update dg-error. * gfortran.dg/gomp/teams-5.f90: New test. gcc/fortran/gfortran.h | 1 + gcc/fortran/openmp.cc | 39 ++++++++- gcc/fortran/parse.cc | 33 ++++++++ gcc/testsuite/gfortran.dg/gomp/pr99226.f90 | 2 +- gcc/testsuite/gfortran.dg/gomp/teams-5.f90 | 127 +++++++++++++++++++++++++++++ 5 files changed, 200 insertions(+), 2 deletions(-) diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h index 6482a885211..577ef807af7 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -1575,6 +1575,7 @@ typedef struct gfc_omp_clauses unsigned order_unconstrained:1, order_reproducible:1, capture:1; unsigned grainsize_strict:1, num_tasks_strict:1, compare:1, weak:1; unsigned non_rectangular:1, order_concurrent:1; + unsigned contains_teams_construct:1; ENUM_BITFIELD (gfc_omp_sched_kind) sched_kind:3; ENUM_BITFIELD (gfc_omp_device_type) device_type:2; ENUM_BITFIELD (gfc_omp_memorder) memorder:3; diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc index 05a697da071..675011a18ce 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -10653,6 +10653,41 @@ gfc_resolve_oacc_directive (gfc_code *code, gfc_namespace *ns ATTRIBUTE_UNUSED) } +static void +resolve_omp_target (gfc_code *code) +{ +#define GFC_IS_TEAMS_CONSTRUCT(op) \ + (op == EXEC_OMP_TEAMS \ + || op == EXEC_OMP_TEAMS_DISTRIBUTE \ + || op == EXEC_OMP_TEAMS_DISTRIBUTE_SIMD \ + || op == EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO \ + || op == EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD \ + || op == EXEC_OMP_TEAMS_LOOP) + + if (!code->ext.omp_clauses->contains_teams_construct) + return; + if ((GFC_IS_TEAMS_CONSTRUCT (code->block->next->op) + && code->block->next->next == NULL) + || (code->block->next->op == EXEC_BLOCK + && code->block->next->next + && GFC_IS_TEAMS_CONSTRUCT (code->block->next->next->op) + && code->block->next->next->next == NULL)) + return; + gfc_code *c = code->block->next; + while (c && !GFC_IS_TEAMS_CONSTRUCT (c->op)) + c = c->next; + if (c) + gfc_error ("!$OMP TARGET region at %L with a nested TEAMS at %L may not " + "contain any other statement, declaration or directive outside " + "of the single TEAMS construct", &c->loc, &code->loc); + else + gfc_error ("!$OMP TARGET region at %L with a nested TEAMS may not " + "contain any other statement, declaration or directive outside " + "of the single TEAMS construct", &code->loc); +#undef GFC_IS_TEAMS_CONSTRUCT +} + + /* Resolve OpenMP directive clauses and check various requirements of each directive. */ @@ -10703,6 +10738,9 @@ gfc_resolve_omp_directive (gfc_code *code, gfc_namespace *ns) case EXEC_OMP_TEAMS_LOOP: resolve_omp_do (code); break; + case EXEC_OMP_TARGET: + resolve_omp_target (code); + gcc_fallthrough (); case EXEC_OMP_ALLOCATE: case EXEC_OMP_ALLOCATORS: case EXEC_OMP_ASSUME: @@ -10718,7 +10756,6 @@ gfc_resolve_omp_directive (gfc_code *code, gfc_namespace *ns) case EXEC_OMP_SCOPE: case EXEC_OMP_SECTIONS: case EXEC_OMP_SINGLE: - case EXEC_OMP_TARGET: case EXEC_OMP_TARGET_DATA: case EXEC_OMP_TARGET_ENTER_DATA: case EXEC_OMP_TARGET_EXIT_DATA: diff --git a/gcc/fortran/parse.cc b/gcc/fortran/parse.cc index e53b7a42e92..011a39c3d04 100644 --- a/gcc/fortran/parse.cc +++ b/gcc/fortran/parse.cc @@ -1312,6 +1312,39 @@ decode_omp_directive (void) prog_unit->omp_target_seen = true; break; } + case ST_OMP_TEAMS: + case ST_OMP_TEAMS_DISTRIBUTE: + case ST_OMP_TEAMS_DISTRIBUTE_SIMD: + case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO: + case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD: + case ST_OMP_TEAMS_LOOP: + if (gfc_state_stack->previous && gfc_state_stack->previous->tail) + { + gfc_state_data *stk = gfc_state_stack; + do { + stk = stk->previous; + } while (stk && stk->tail && stk->tail->op == EXEC_BLOCK); + if (stk && stk->tail) + switch (stk->tail->op) + { + case EXEC_OMP_TARGET: + case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE: + case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD: + case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO: + case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD: + case EXEC_OMP_TARGET_TEAMS_LOOP: + case EXEC_OMP_TARGET_PARALLEL: + case EXEC_OMP_TARGET_PARALLEL_DO: + case EXEC_OMP_TARGET_PARALLEL_DO_SIMD: + case EXEC_OMP_TARGET_PARALLEL_LOOP: + case EXEC_OMP_TARGET_SIMD: + stk->tail->ext.omp_clauses->contains_teams_construct = 1; + break; + default: + break; + } + } + break; case ST_OMP_ERROR: if (new_st.ext.omp_clauses->at != OMP_AT_EXECUTION) return ST_NONE; diff --git a/gcc/testsuite/gfortran.dg/gomp/pr99226.f90 b/gcc/testsuite/gfortran.dg/gomp/pr99226.f90 index 72dbdde2e28..2aea0c15585 100644 --- a/gcc/testsuite/gfortran.dg/gomp/pr99226.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/pr99226.f90 @@ -2,7 +2,7 @@ subroutine sub (n) integer :: n, i - !$omp target ! { dg-error "construct with nested 'teams' construct contains directives outside of the 'teams' construct" } + !$omp target ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } !$omp teams distribute dist_schedule (static,n+4) do i = 1, 8 end do diff --git a/gcc/testsuite/gfortran.dg/gomp/teams-5.f90 b/gcc/testsuite/gfortran.dg/gomp/teams-5.f90 new file mode 100644 index 00000000000..bf5461b87c8 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/teams-5.f90 @@ -0,0 +1,127 @@ +! { dg-do compile } + +! PR fortran/110725 +! PR middle-end/71065 + +implicit none +integer :: x +!$omp target device(1) + block + !$omp teams num_teams(f()) + !$omp end teams + end block +!!$omp end target + +!$omp target device(1) + !$omp teams num_teams(f()) + !$omp end teams +!$omp end target + +!$omp target device(1) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + x = 5 + !$omp teams num_teams(f()) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + !$omp end teams +!$omp end target + +!$omp target device(1) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + !$omp teams num_teams(f()) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + !$omp end teams + x = 5 +!$omp end target + +!$omp target ! { dg-error "OMP TARGET region at .1. with a nested TEAMS may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + block + block + !$omp teams num_teams(f()) + !$omp end teams + end block + end block +!$omp end target + +!$omp target device(1) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + block + x = 5 + !$omp teams num_teams(f()) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + !$omp end teams + end block +!$omp end target + +!$omp target device(1) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + block + !$omp teams num_teams(f()) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + !$omp end teams + x = 5 + end block +!$omp end target + +!$omp target device(1) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + block + !$omp teams num_teams(f()) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + !$omp end teams + x = 5 + end block +!$omp end target + +contains + +function f() + !$omp declare target + integer, allocatable :: f + f = 5 +end +end + +subroutine sub1 + implicit none + integer :: x,i + + !$omp target device(1) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + block + !$omp teams distribute num_teams(f()) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + do i = 1, 5 + end do + x = 7 + end block + !$omp end target + + !$omp target device(1) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + block + !$omp teams loop num_teams(f()) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + do i = 1, 5 + end do + x = 7 + end block + !$omp end target + + !$omp target device(1) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + !$omp teams distribute simd num_teams(f()) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + do i = 1, 5 + end do + x = 7 + !$omp end target + + !$omp target device(1) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + !$omp teams distribute parallel do num_teams(f()) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + do i = 1, 5 + end do + x = 7 + !$omp end target + + !$omp target device(1) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + block + x = 7 + !$omp teams distribute parallel do simd num_teams(f()) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + do i = 1, 5 + end do + end block + !$omp end target + +contains + +function f() + !$omp declare target + integer, allocatable :: f + f = 5 +end + +end ^ permalink raw reply [flat|nested] 6+ messages in thread
* Re: [Patch] OpenMP/Fortran: Reject not strictly nested target -> teams [PR110725, PR71065] 2023-07-24 19:43 [Patch] OpenMP/Fortran: Reject not strictly nested target -> teams [PR110725, PR71065] Tobias Burnus @ 2023-07-24 19:49 ` Jakub Jelinek 2023-07-24 20:05 ` Tobias Burnus 2023-07-25 11:14 ` [patch] OpenMP/Fortran: Reject declarations between target + teams (was: [Patch] OpenMP/Fortran: Reject not strictly nested target -> teams [PR110725, PR71065]) Tobias Burnus 0 siblings, 2 replies; 6+ messages in thread From: Jakub Jelinek @ 2023-07-24 19:49 UTC (permalink / raw) To: Tobias Burnus; +Cc: gcc-patches, fortran On Mon, Jul 24, 2023 at 09:43:10PM +0200, Tobias Burnus wrote: > This patch adds diagnostic for additional code alongside a nested teams > in a target region. > > The diagnostic is happening soon after parsing such that expressions > in clauses are not yet expanded - those would end up before TEAMS > and can be very complicated (e.g. assume an allocatable-returning function). > > (The patch diagnoses it in openmp.cc; after trans-openmp.cc it would > already be to late.) > > Comments, remarks, suggestions? Thanks for working on this. The fuzzy thing on the Fortran side is if e.g. multiple nested BLOCK statements can appear sandwiched in between target and teams (of course without declarations in them), or if e.g. extra empty BLOCK; END BLOCK could appear next to it etc. And on C/C++ side similarly with {}s, ; is an empty statement, so #pragma omp target { ; #pragma omp teams ; ; } etc. would be invalid. Jakub ^ permalink raw reply [flat|nested] 6+ messages in thread
* Re: [Patch] OpenMP/Fortran: Reject not strictly nested target -> teams [PR110725, PR71065] 2023-07-24 19:49 ` Jakub Jelinek @ 2023-07-24 20:05 ` Tobias Burnus 2023-07-25 7:37 ` Tobias Burnus 2023-07-25 11:14 ` [patch] OpenMP/Fortran: Reject declarations between target + teams (was: [Patch] OpenMP/Fortran: Reject not strictly nested target -> teams [PR110725, PR71065]) Tobias Burnus 1 sibling, 1 reply; 6+ messages in thread From: Tobias Burnus @ 2023-07-24 20:05 UTC (permalink / raw) To: Jakub Jelinek; +Cc: gcc-patches, fortran On 24.07.23 21:49, Jakub Jelinek via Fortran wrote: > Thanks for working on this. The fuzzy thing on the Fortran side is > if e.g. multiple nested BLOCK statements can appear sandwiched in between > target and teams (of course without declarations in them), or if e.g. The current patch rejects nested blocks, be it 'omp target; block; block; omp teams;' or be it 'omp target; block; block;end block; omp teams'. The current wording in the spec is also rather explicit as 'block' is a statement. (BTW: For 'block; block; omp teams', the simplistic search won't work such that for those only the location of TARGET and not of TEAMS is shown. I could try harder but as it is useful as is and such code should be rare, I don't do it.) Thus, I believe the patch should be fine. Tobias PS: I know that some regard {{{ }}} and block; block, ... end block; ... as something to be ignored. Thus, for 'omp atomic', TR12 will allow any number of curly braces and BLOCK/ENDBLOCK pairs. The wording there is rather explicit but also localized, i.e. it won't affect other code locations (for now at least). ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 ^ permalink raw reply [flat|nested] 6+ messages in thread
* Re: [Patch] OpenMP/Fortran: Reject not strictly nested target -> teams [PR110725, PR71065] 2023-07-24 20:05 ` Tobias Burnus @ 2023-07-25 7:37 ` Tobias Burnus 0 siblings, 0 replies; 6+ messages in thread From: Tobias Burnus @ 2023-07-25 7:37 UTC (permalink / raw) To: gcc-patches; +Cc: fortran, Jakub Jelinek [-- Attachment #1: Type: text/plain, Size: 744 bytes --] Now committed as r14-2754-g2e31fe431b08b0 with a minor addition: On 24.07.23 22:05, Tobias Burnus wrote: > The current patch rejects nested blocks, be it 'omp target; block; > block; omp teams;' which was before in the testcase. But now also > or be it 'omp target; block; block;end block; omp teams'. is tested for. Somehow, the second dg-error line in an modified testcase did not make it in the first commit; now fixed in r14-2759-g50656980497d77 Tobias. ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 [-- Attachment #2: committed.diff --] [-- Type: text/x-patch, Size: 1657 bytes --] commit 50656980497d77ac12a5e7179013a6af09ba32f7 Author: Tobias Burnus <tobias@codesourcery.com> Date: Tue Jul 25 09:28:43 2023 +0200 gfortran.dg/gomp/pr99226.f90: Add missing dg-error Follow up to r14-2754-g2e31fe431b08b0302e1fa8a1c18ee51adafd41df which added a check that a target region with teams does not have anything anything else strictly nested in the target. When changing the dg-error for this PR, somehow the addition of a dg-error in a second line was lost (the message uses (1) and (2) as location, showing two lines, both need a dg-error with the same message). gcc/testsuite/ChangeLog: * gfortran.dg/gomp/pr99226.f90: Update dg-error. --- gcc/testsuite/gfortran.dg/gomp/pr99226.f90 | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/gcc/testsuite/gfortran.dg/gomp/pr99226.f90 b/gcc/testsuite/gfortran.dg/gomp/pr99226.f90 index 2aea0c15585..d1b35076dd0 100644 --- a/gcc/testsuite/gfortran.dg/gomp/pr99226.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/pr99226.f90 @@ -3,7 +3,7 @@ subroutine sub (n) integer :: n, i !$omp target ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } - !$omp teams distribute dist_schedule (static,n+4) + !$omp teams distribute dist_schedule (static,n+4) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } do i = 1, 8 end do !$omp teams distribute dist_schedule (static,n+4) [-- Attachment #3: committed2.diff --] [-- Type: text/x-patch, Size: 13783 bytes --] commit 2e31fe431b08b0302e1fa8a1c18ee51adafd41df Author: Tobias Burnus <tobias@codesourcery.com> Date: Mon Jul 24 22:57:07 2023 +0200 OpenMP/Fortran: Reject not strictly nested target -> teams [PR110725, PR71065] OpenMP requires: "If a teams region is nested inside a target region, the corresponding target construct must not contain any statements, declarations or directives outside of the corresponding teams construct." This commit checks now for this restriction. PR fortran/110725 PR middle-end/71065 gcc/fortran/ChangeLog: * gfortran.h (gfc_omp_clauses): Add contains_teams_construct. * openmp.cc (resolve_omp_target): New; check for teams nesting. (gfc_resolve_omp_directive): Call it. * parse.cc (decode_omp_directive): Set contains_teams_construct on enclosing ST_OMP_TARGET. gcc/testsuite/ChangeLog: * gfortran.dg/gomp/pr99226.f90: Update dg-error. * gfortran.dg/gomp/teams-5.f90: New test. --- gcc/fortran/gfortran.h | 1 + gcc/fortran/openmp.cc | 39 +++++++- gcc/fortran/parse.cc | 33 +++++++ gcc/testsuite/gfortran.dg/gomp/pr99226.f90 | 2 +- gcc/testsuite/gfortran.dg/gomp/teams-5.f90 | 150 +++++++++++++++++++++++++++++ 5 files changed, 223 insertions(+), 2 deletions(-) diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h index 6482a885211..577ef807af7 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -1575,6 +1575,7 @@ typedef struct gfc_omp_clauses unsigned order_unconstrained:1, order_reproducible:1, capture:1; unsigned grainsize_strict:1, num_tasks_strict:1, compare:1, weak:1; unsigned non_rectangular:1, order_concurrent:1; + unsigned contains_teams_construct:1; ENUM_BITFIELD (gfc_omp_sched_kind) sched_kind:3; ENUM_BITFIELD (gfc_omp_device_type) device_type:2; ENUM_BITFIELD (gfc_omp_memorder) memorder:3; diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc index 05a697da071..675011a18ce 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -10653,6 +10653,41 @@ gfc_resolve_oacc_directive (gfc_code *code, gfc_namespace *ns ATTRIBUTE_UNUSED) } +static void +resolve_omp_target (gfc_code *code) +{ +#define GFC_IS_TEAMS_CONSTRUCT(op) \ + (op == EXEC_OMP_TEAMS \ + || op == EXEC_OMP_TEAMS_DISTRIBUTE \ + || op == EXEC_OMP_TEAMS_DISTRIBUTE_SIMD \ + || op == EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO \ + || op == EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD \ + || op == EXEC_OMP_TEAMS_LOOP) + + if (!code->ext.omp_clauses->contains_teams_construct) + return; + if ((GFC_IS_TEAMS_CONSTRUCT (code->block->next->op) + && code->block->next->next == NULL) + || (code->block->next->op == EXEC_BLOCK + && code->block->next->next + && GFC_IS_TEAMS_CONSTRUCT (code->block->next->next->op) + && code->block->next->next->next == NULL)) + return; + gfc_code *c = code->block->next; + while (c && !GFC_IS_TEAMS_CONSTRUCT (c->op)) + c = c->next; + if (c) + gfc_error ("!$OMP TARGET region at %L with a nested TEAMS at %L may not " + "contain any other statement, declaration or directive outside " + "of the single TEAMS construct", &c->loc, &code->loc); + else + gfc_error ("!$OMP TARGET region at %L with a nested TEAMS may not " + "contain any other statement, declaration or directive outside " + "of the single TEAMS construct", &code->loc); +#undef GFC_IS_TEAMS_CONSTRUCT +} + + /* Resolve OpenMP directive clauses and check various requirements of each directive. */ @@ -10703,6 +10738,9 @@ gfc_resolve_omp_directive (gfc_code *code, gfc_namespace *ns) case EXEC_OMP_TEAMS_LOOP: resolve_omp_do (code); break; + case EXEC_OMP_TARGET: + resolve_omp_target (code); + gcc_fallthrough (); case EXEC_OMP_ALLOCATE: case EXEC_OMP_ALLOCATORS: case EXEC_OMP_ASSUME: @@ -10718,7 +10756,6 @@ gfc_resolve_omp_directive (gfc_code *code, gfc_namespace *ns) case EXEC_OMP_SCOPE: case EXEC_OMP_SECTIONS: case EXEC_OMP_SINGLE: - case EXEC_OMP_TARGET: case EXEC_OMP_TARGET_DATA: case EXEC_OMP_TARGET_ENTER_DATA: case EXEC_OMP_TARGET_EXIT_DATA: diff --git a/gcc/fortran/parse.cc b/gcc/fortran/parse.cc index e53b7a42e92..011a39c3d04 100644 --- a/gcc/fortran/parse.cc +++ b/gcc/fortran/parse.cc @@ -1312,6 +1312,39 @@ decode_omp_directive (void) prog_unit->omp_target_seen = true; break; } + case ST_OMP_TEAMS: + case ST_OMP_TEAMS_DISTRIBUTE: + case ST_OMP_TEAMS_DISTRIBUTE_SIMD: + case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO: + case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD: + case ST_OMP_TEAMS_LOOP: + if (gfc_state_stack->previous && gfc_state_stack->previous->tail) + { + gfc_state_data *stk = gfc_state_stack; + do { + stk = stk->previous; + } while (stk && stk->tail && stk->tail->op == EXEC_BLOCK); + if (stk && stk->tail) + switch (stk->tail->op) + { + case EXEC_OMP_TARGET: + case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE: + case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD: + case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO: + case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD: + case EXEC_OMP_TARGET_TEAMS_LOOP: + case EXEC_OMP_TARGET_PARALLEL: + case EXEC_OMP_TARGET_PARALLEL_DO: + case EXEC_OMP_TARGET_PARALLEL_DO_SIMD: + case EXEC_OMP_TARGET_PARALLEL_LOOP: + case EXEC_OMP_TARGET_SIMD: + stk->tail->ext.omp_clauses->contains_teams_construct = 1; + break; + default: + break; + } + } + break; case ST_OMP_ERROR: if (new_st.ext.omp_clauses->at != OMP_AT_EXECUTION) return ST_NONE; diff --git a/gcc/testsuite/gfortran.dg/gomp/pr99226.f90 b/gcc/testsuite/gfortran.dg/gomp/pr99226.f90 index 72dbdde2e28..2aea0c15585 100644 --- a/gcc/testsuite/gfortran.dg/gomp/pr99226.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/pr99226.f90 @@ -2,7 +2,7 @@ subroutine sub (n) integer :: n, i - !$omp target ! { dg-error "construct with nested 'teams' construct contains directives outside of the 'teams' construct" } + !$omp target ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } !$omp teams distribute dist_schedule (static,n+4) do i = 1, 8 end do diff --git a/gcc/testsuite/gfortran.dg/gomp/teams-5.f90 b/gcc/testsuite/gfortran.dg/gomp/teams-5.f90 new file mode 100644 index 00000000000..00377b69bf4 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/teams-5.f90 @@ -0,0 +1,150 @@ +! { dg-do compile } + +! PR fortran/110725 +! PR middle-end/71065 + +implicit none +integer :: x +!$omp target device(1) + block + !$omp teams num_teams(f()) + !$omp end teams + end block +!!$omp end target + +!$omp target device(1) + !$omp teams num_teams(f()) + !$omp end teams +!$omp end target + +!$omp target device(1) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + x = 5 + !$omp teams num_teams(f()) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + !$omp end teams +!$omp end target + +!$omp target device(1) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + !$omp teams num_teams(f()) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + !$omp end teams + x = 5 +!$omp end target + +!$omp target ! { dg-error "OMP TARGET region at .1. with a nested TEAMS may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + block + block + !$omp teams num_teams(f()) + !$omp end teams + end block + end block +!$omp end target + +!$omp target device(1) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + block + x = 5 + !$omp teams num_teams(f()) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + !$omp end teams + end block +!$omp end target + +!$omp target device(1) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + block + !$omp teams num_teams(f()) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + !$omp end teams + x = 5 + end block +!$omp end target + +!$omp target device(1) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + block + !$omp teams num_teams(f()) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + !$omp end teams + x = 5 + end block +!$omp end target + +!$omp target device(1) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + !$omp teams num_teams(f()) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + !$omp end teams + block; end block +!$omp end target + +!$omp target device(1) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + block + block; end block; + !$omp teams num_teams(f()) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + !$omp end teams + end block +!$omp end target + +!$omp target device(1) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + block + !$omp teams num_teams(f()) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + !$omp end teams + block; end block; + end block +!!$omp end target + + +contains + +function f() + !$omp declare target + integer, allocatable :: f + f = 5 +end +end + +subroutine sub1 + implicit none + integer :: x,i + + !$omp target device(1) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + block + !$omp teams distribute num_teams(f()) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + do i = 1, 5 + end do + x = 7 + end block + !$omp end target + + !$omp target device(1) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + block + !$omp teams loop num_teams(f()) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + do i = 1, 5 + end do + x = 7 + end block + !$omp end target + + !$omp target device(1) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + !$omp teams distribute simd num_teams(f()) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + do i = 1, 5 + end do + x = 7 + !$omp end target + + !$omp target device(1) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + !$omp teams distribute parallel do num_teams(f()) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + do i = 1, 5 + end do + x = 7 + !$omp end target + + !$omp target device(1) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + block + x = 7 + !$omp teams distribute parallel do simd num_teams(f()) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + do i = 1, 5 + end do + end block + !$omp end target + +contains + +function f() + !$omp declare target + integer, allocatable :: f + f = 5 +end + +end ^ permalink raw reply [flat|nested] 6+ messages in thread
* [patch] OpenMP/Fortran: Reject declarations between target + teams (was: [Patch] OpenMP/Fortran: Reject not strictly nested target -> teams [PR110725, PR71065]) 2023-07-24 19:49 ` Jakub Jelinek 2023-07-24 20:05 ` Tobias Burnus @ 2023-07-25 11:14 ` Tobias Burnus 2023-07-27 16:36 ` [committed] OpenMP/Fortran: Extend reject code between target + teams [PR71065, PR110725] (was: Re: [patch] OpenMP/Fortran: Reject declarations between target + teams (was: [Patch] OpenMP/Fortran: Reject not strictly nested target -> teams [PR110725, PR71065])) Tobias Burnus 1 sibling, 1 reply; 6+ messages in thread From: Tobias Burnus @ 2023-07-25 11:14 UTC (permalink / raw) To: Jakub Jelinek, gcc-patches; +Cc: fortran [-- Attachment #1: Type: text/plain, Size: 1021 bytes --] On 24.07.23 21:49, Jakub Jelinek via Fortran wrote: > On Mon, Jul 24, 2023 at 09:43:10PM +0200, Tobias Burnus wrote: >> This patch adds diagnostic for additional code alongside a nested teams >> in a target region. > Thanks for working on this. The fuzzy thing on the Fortran side is > if e.g. multiple nested BLOCK statements can appear sandwiched in between > target and teams (of course without declarations in them) Talking about declarations, I realized that I missed to diagnose them; the attached patch should handle them as well. (Except for 'omp nothing' and 'omp error', which return ST_NONE.) Comments, remarks, suggestions? If none or no changes are required, I will later commit the attached follow-up patch. Tobias ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 [-- Attachment #2: omp-target-teams-diag-decl.diff --] [-- Type: text/x-patch, Size: 7882 bytes --] OpenMP/Fortran: Reject declarations between target + teams While commit r14-2754-g2e31fe431b08b0302e1fa8a1c18ee51adafd41df detected executable statements, declarations do not show up as executable statements. Hence, we now check whether the first statement after TARGET is TEAMS - such that we can detect data statements like type or variable declarations. Fortran semantics ensures that only executable directives/statemens can come after '!$omp end teams' such that those can be detected with the previous check. Note that statements returning ST_NONE such as 'omp nothing' or 'omp error at(compilation)' will still slip through. PR fortran/110725 PR middle-end/71065 gcc/fortran/ChangeLog: * gfortran.h (gfc_omp_clauses): Add target_first_st_is_teams. * parse.cc (parse_omp_structured_block): Set it if the first statement in the structured block of a TARGET is TEAMS or a combined/composite starting with TEAMS. * openmp.cc (resolve_omp_target): Also show an error for contains_teams_construct without target_first_st_is_teams. gcc/testsuite/ChangeLog: * gfortran.dg/gomp/teams-6.f90: New test. gcc/fortran/gfortran.h | 2 +- gcc/fortran/openmp.cc | 13 ++--- gcc/fortran/parse.cc | 25 ++++++++-- gcc/testsuite/gfortran.dg/gomp/teams-6.f90 | 78 ++++++++++++++++++++++++++++++ 4 files changed, 108 insertions(+), 10 deletions(-) diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h index 577ef807af7..9a00e6dea6f 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -1575,7 +1575,7 @@ typedef struct gfc_omp_clauses unsigned order_unconstrained:1, order_reproducible:1, capture:1; unsigned grainsize_strict:1, num_tasks_strict:1, compare:1, weak:1; unsigned non_rectangular:1, order_concurrent:1; - unsigned contains_teams_construct:1; + unsigned contains_teams_construct:1, target_first_st_is_teams:1; ENUM_BITFIELD (gfc_omp_sched_kind) sched_kind:3; ENUM_BITFIELD (gfc_omp_device_type) device_type:2; ENUM_BITFIELD (gfc_omp_memorder) memorder:3; diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc index 675011a18ce..52eeaf2d4da 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -10666,12 +10666,13 @@ resolve_omp_target (gfc_code *code) if (!code->ext.omp_clauses->contains_teams_construct) return; - if ((GFC_IS_TEAMS_CONSTRUCT (code->block->next->op) - && code->block->next->next == NULL) - || (code->block->next->op == EXEC_BLOCK - && code->block->next->next - && GFC_IS_TEAMS_CONSTRUCT (code->block->next->next->op) - && code->block->next->next->next == NULL)) + if (code->ext.omp_clauses->target_first_st_is_teams + && ((GFC_IS_TEAMS_CONSTRUCT (code->block->next->op) + && code->block->next->next == NULL) + || (code->block->next->op == EXEC_BLOCK + && code->block->next->next + && GFC_IS_TEAMS_CONSTRUCT (code->block->next->next->op) + && code->block->next->next->next == NULL))) return; gfc_code *c = code->block->next; while (c && !GFC_IS_TEAMS_CONSTRUCT (c->op)) diff --git a/gcc/fortran/parse.cc b/gcc/fortran/parse.cc index 011a39c3d04..aa6bb663def 100644 --- a/gcc/fortran/parse.cc +++ b/gcc/fortran/parse.cc @@ -5766,7 +5766,7 @@ parse_openmp_allocate_block (gfc_statement omp_st) static gfc_statement parse_omp_structured_block (gfc_statement omp_st, bool workshare_stmts_only) { - gfc_statement st, omp_end_st; + gfc_statement st, omp_end_st, first_st; gfc_code *cp, *np; gfc_state_data s; @@ -5857,7 +5857,7 @@ parse_omp_structured_block (gfc_statement omp_st, bool workshare_stmts_only) gfc_namespace *my_ns = NULL; gfc_namespace *my_parent = NULL; - st = next_statement (); + first_st = st = next_statement (); if (st == ST_BLOCK) { @@ -5876,9 +5876,28 @@ parse_omp_structured_block (gfc_statement omp_st, bool workshare_stmts_only) new_st.ext.block.ns = my_ns; new_st.ext.block.assoc = NULL; accept_statement (ST_BLOCK); - st = parse_spec (ST_NONE); + first_st = next_statement (); + st = parse_spec (first_st); } + if (omp_end_st == ST_OMP_END_TARGET) + switch (first_st) + { + case ST_OMP_TEAMS: + case ST_OMP_TEAMS_DISTRIBUTE: + case ST_OMP_TEAMS_DISTRIBUTE_SIMD: + case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO: + case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD: + case ST_OMP_TEAMS_LOOP: + { + gfc_state_data *stk = gfc_state_stack->previous; + stk->tail->ext.omp_clauses->target_first_st_is_teams = true; + break; + } + default: + break; + } + do { if (workshare_stmts_only) diff --git a/gcc/testsuite/gfortran.dg/gomp/teams-6.f90 b/gcc/testsuite/gfortran.dg/gomp/teams-6.f90 new file mode 100644 index 00000000000..be453f27f40 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/teams-6.f90 @@ -0,0 +1,78 @@ +! { dg-do compile } + +! PR fortran/110725 +! PR middle-end/71065 + + +subroutine one +!$omp target ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } +block + integer :: i ! <<< invalid: variable declaration + !$omp teams ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + i = 5 + !$omp end teams +end block + +!$omp target ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } +block + type t ! <<< invalid: type declaration + end type t + !$omp teams ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + i = 5 + !$omp end teams +end block + +!$omp target + ! The following is invalid - but not detected as ST_NONE is returned: + !$omp error at(compilation) severity(warning) ! { dg-warning "OMP ERROR encountered" } + !$omp teams + i = 5 + !$omp end teams +!$omp end target + +!$omp target + ! The following is invalid - but not detected as ST_NONE is returned: + !$omp nothing ! <<< invalid: directive + !$omp teams + i = 5 + !$omp end teams +!$omp end target +end + + +subroutine two +!$omp target ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } +block + integer :: i ! <<< invalid: variable declaration + !$omp teams distribute ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + do i = 1, 5 + end do + !$omp end teams distribute +end block + +!$omp target ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } +block + type t ! <<< invalid: type declaration + end type t + !$omp teams distribute parallel do ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + do i = 1, 5 + end do +end block + +!$omp target + ! The following is invalid - but not detected as ST_NONE is returned: + !$omp error at(compilation) severity(warning) ! { dg-warning "OMP ERROR encountered" } + !$omp teams loop + do i = 5, 10 + end do +!$omp end target + +!$omp target + ! The following is invalid - but not detected as ST_NONE is returned: + !$omp nothing ! <<< invalid: directive + !$omp teams distribute simd + do i = -3, 5 + end do + !$omp end teams distribute simd +!$omp end target +end ^ permalink raw reply [flat|nested] 6+ messages in thread
* [committed] OpenMP/Fortran: Extend reject code between target + teams [PR71065, PR110725] (was: Re: [patch] OpenMP/Fortran: Reject declarations between target + teams (was: [Patch] OpenMP/Fortran: Reject not strictly nested target -> teams [PR110725, PR71065])) 2023-07-25 11:14 ` [patch] OpenMP/Fortran: Reject declarations between target + teams (was: [Patch] OpenMP/Fortran: Reject not strictly nested target -> teams [PR110725, PR71065]) Tobias Burnus @ 2023-07-27 16:36 ` Tobias Burnus 0 siblings, 0 replies; 6+ messages in thread From: Tobias Burnus @ 2023-07-27 16:36 UTC (permalink / raw) To: gcc-patches; +Cc: fortran, Jakub Jelinek [-- Attachment #1: Type: text/plain, Size: 658 bytes --] Yet another omission, the flag was not properly set for deeply buried 'omp teams' as I stopped too early when walking up the stack. Now fixed by commit r14-2826-g081e25d3cfd86c * * * This was found when 'repairing' the feature on the OG13 (devel/omp/gcc-13) branch for metadirectives, cf. the second attached patch, applied after cherry-picking the mainline patch. Tobias ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 [-- Attachment #2: committed.diff --] [-- Type: text/x-patch, Size: 4414 bytes --] commit 081e25d3cfd86c4094999ded0bbe99b91762013c Author: Tobias Burnus <tobias@codesourcery.com> Date: Thu Jul 27 18:14:11 2023 +0200 OpenMP/Fortran: Extend reject code between target + teams [PR71065, PR110725] The previous version failed to diagnose when the 'teams' was nested more deeply inside the target region, e.g. inside a DO or some block or structured block. PR fortran/110725 PR middle-end/71065 gcc/fortran/ChangeLog: * openmp.cc (resolve_omp_target): Minor cleanup. * parse.cc (decode_omp_directive): Find TARGET statement also higher in the stack. gcc/testsuite/ChangeLog: * gfortran.dg/gomp/teams-6.f90: Extend. diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc index 52eeaf2d4da..2952cd300ac 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -10666,15 +10666,14 @@ resolve_omp_target (gfc_code *code) if (!code->ext.omp_clauses->contains_teams_construct) return; + gfc_code *c = code->block->next; if (code->ext.omp_clauses->target_first_st_is_teams - && ((GFC_IS_TEAMS_CONSTRUCT (code->block->next->op) - && code->block->next->next == NULL) - || (code->block->next->op == EXEC_BLOCK - && code->block->next->next - && GFC_IS_TEAMS_CONSTRUCT (code->block->next->next->op) - && code->block->next->next->next == NULL))) + && ((GFC_IS_TEAMS_CONSTRUCT (c->op) && c->next == NULL) + || (c->op == EXEC_BLOCK + && c->next + && GFC_IS_TEAMS_CONSTRUCT (c->next->op) + && c->next->next == NULL))) return; - gfc_code *c = code->block->next; while (c && !GFC_IS_TEAMS_CONSTRUCT (c->op)) c = c->next; if (c) diff --git a/gcc/fortran/parse.cc b/gcc/fortran/parse.cc index aa6bb663def..e797402b59f 100644 --- a/gcc/fortran/parse.cc +++ b/gcc/fortran/parse.cc @@ -1318,32 +1318,27 @@ decode_omp_directive (void) case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO: case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD: case ST_OMP_TEAMS_LOOP: - if (gfc_state_stack->previous && gfc_state_stack->previous->tail) - { - gfc_state_data *stk = gfc_state_stack; - do { - stk = stk->previous; - } while (stk && stk->tail && stk->tail->op == EXEC_BLOCK); - if (stk && stk->tail) - switch (stk->tail->op) - { - case EXEC_OMP_TARGET: - case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE: - case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD: - case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO: - case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD: - case EXEC_OMP_TARGET_TEAMS_LOOP: - case EXEC_OMP_TARGET_PARALLEL: - case EXEC_OMP_TARGET_PARALLEL_DO: - case EXEC_OMP_TARGET_PARALLEL_DO_SIMD: - case EXEC_OMP_TARGET_PARALLEL_LOOP: - case EXEC_OMP_TARGET_SIMD: - stk->tail->ext.omp_clauses->contains_teams_construct = 1; - break; - default: - break; - } - } + for (gfc_state_data *stk = gfc_state_stack->previous; stk; + stk = stk->previous) + if (stk && stk->tail) + switch (stk->tail->op) + { + case EXEC_OMP_TARGET: + case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE: + case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD: + case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO: + case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD: + case EXEC_OMP_TARGET_TEAMS_LOOP: + case EXEC_OMP_TARGET_PARALLEL: + case EXEC_OMP_TARGET_PARALLEL_DO: + case EXEC_OMP_TARGET_PARALLEL_DO_SIMD: + case EXEC_OMP_TARGET_PARALLEL_LOOP: + case EXEC_OMP_TARGET_SIMD: + stk->tail->ext.omp_clauses->contains_teams_construct = 1; + break; + default: + break; + } break; case ST_OMP_ERROR: if (new_st.ext.omp_clauses->at != OMP_AT_EXECUTION) diff --git a/gcc/testsuite/gfortran.dg/gomp/teams-6.f90 b/gcc/testsuite/gfortran.dg/gomp/teams-6.f90 index be453f27f40..0bd7735e738 100644 --- a/gcc/testsuite/gfortran.dg/gomp/teams-6.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/teams-6.f90 @@ -37,6 +37,16 @@ end block i = 5 !$omp end teams !$omp end target + + +!$omp target ! { dg-error "OMP TARGET region at .1. with a nested TEAMS may not contain any other statement, declaration or directive outside of the single TEAMS construct" } +block + do i = 5, 8 + !$omp teams + block; end block + end do +end block + end [-- Attachment #3: og13.diff --] [-- Type: text/x-patch, Size: 8489 bytes --] commit eae457d9aa6ccad1692759bffee8fa3f6c92a3a0 Author: Tobias Burnus <tobias@codesourcery.com> Date: Thu Jul 27 18:30:20 2023 +0200 OpenMP/Fortran: Fix target + teams diagnostic with metadirectives gcc/fortran/ChangeLog: * gfortran.h (gfc_omp_clauses): Rename target_first_st_is_teams to target_first_st_is_teams_or_meta. * parse.cc (parse_omp_structured_block): Handle metadirectives for target_first_st_is_teams. * openmp.cc (resolve_omp_target): Likewise to fix target+teams diagnostic with metadirectives. libgomp/ChangeLog: * testsuite/libgomp.fortran/metadirective-1.f90: Extend. * testsuite/libgomp.fortran/metadirective-6.f90: New test. --- gcc/fortran/ChangeLog.omp | 9 ++++ gcc/fortran/gfortran.h | 2 +- gcc/fortran/openmp.cc | 35 ++++++++++--- gcc/fortran/parse.cc | 4 +- libgomp/ChangeLog.omp | 5 ++ .../testsuite/libgomp.fortran/metadirective-1.f90 | 28 +++++++++++ .../testsuite/libgomp.fortran/metadirective-6.f90 | 58 ++++++++++++++++++++++ 7 files changed, 132 insertions(+), 9 deletions(-) diff --git a/gcc/fortran/ChangeLog.omp b/gcc/fortran/ChangeLog.omp index c197f77f1f9..237e9ebeba2 100644 --- a/gcc/fortran/ChangeLog.omp +++ b/gcc/fortran/ChangeLog.omp @@ -1,3 +1,12 @@ +2023-07-27 Tobias Burnus <tobias@codesourcery.com> + + * gfortran.h (gfc_omp_clauses): Rename target_first_st_is_teams + to target_first_st_is_teams_or_meta. + * parse.cc (parse_omp_structured_block): Handle metadirectives + for target_first_st_is_teams. + * openmp.cc (resolve_omp_target): Likewise to fix target+teams + diagnostic with metadirectives. + 2023-07-27 Tobias Burnus <tobias@codesourcery.com> Backported from master: diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h index 2cf8a0e0c39..0e7e80e4bf1 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -1588,7 +1588,7 @@ typedef struct gfc_omp_clauses unsigned order_unconstrained:1, order_reproducible:1, capture:1; unsigned grainsize_strict:1, num_tasks_strict:1, compare:1, weak:1; unsigned non_rectangular:1, order_concurrent:1; - unsigned contains_teams_construct:1, target_first_st_is_teams:1; + unsigned contains_teams_construct:1, target_first_st_is_teams_or_meta:1; unsigned unroll_full:1, unroll_none:1, unroll_partial:1; unsigned unroll_partial_factor; ENUM_BITFIELD (gfc_omp_sched_kind) sched_kind:3; diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc index 18a4a33feaa..deccb14a525 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -12252,13 +12252,34 @@ resolve_omp_target (gfc_code *code) if (!code->ext.omp_clauses->contains_teams_construct) return; gfc_code *c = code->block->next; - if (code->ext.omp_clauses->target_first_st_is_teams - && ((GFC_IS_TEAMS_CONSTRUCT (c->op) && c->next == NULL) - || (c->op == EXEC_BLOCK - && c->next - && GFC_IS_TEAMS_CONSTRUCT (c->next->op) - && c->next->next == NULL))) - return; + if (c->op == EXEC_BLOCK) + c = c->next; + if (code->ext.omp_clauses->target_first_st_is_teams_or_meta) + { + if (c->op == EXEC_OMP_METADIRECTIVE) + { + struct gfc_omp_metadirective_clause *mc + = c->ext.omp_metadirective_clauses; + /* All mc->(next...->)code should be identical with regards + to the diagnostic below. */ + do + { + if (mc->stmt != ST_NONE + && GFC_IS_TEAMS_CONSTRUCT (mc->code->op)) + { + if (c->next == NULL && mc->code->next == NULL) + return; + c = mc->code; + break; + } + mc = mc->next; + } + while (mc); + } + else if (GFC_IS_TEAMS_CONSTRUCT (c->op) && c->next == NULL) + return; + } + while (c && !GFC_IS_TEAMS_CONSTRUCT (c->op)) c = c->next; if (c) diff --git a/gcc/fortran/parse.cc b/gcc/fortran/parse.cc index 2070a8a7dee..efedde1d84b 100644 --- a/gcc/fortran/parse.cc +++ b/gcc/fortran/parse.cc @@ -5833,9 +5833,11 @@ parse_omp_structured_block (gfc_statement omp_st, bool workshare_stmts_only) case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO: case ST_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD: case ST_OMP_TEAMS_LOOP: + case ST_OMP_METADIRECTIVE: + case ST_OMP_BEGIN_METADIRECTIVE: { gfc_state_data *stk = gfc_state_stack->previous; - stk->tail->ext.omp_clauses->target_first_st_is_teams = true; + stk->tail->ext.omp_clauses->target_first_st_is_teams_or_meta = true; break; } default: diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp index f83700f1c00..9f8e3ec947d 100644 --- a/libgomp/ChangeLog.omp +++ b/libgomp/ChangeLog.omp @@ -1,3 +1,8 @@ +2023-07-27 Tobias Burnus <tobias@codesourcery.com> + + * testsuite/libgomp.fortran/metadirective-1.f90: Extend. + * testsuite/libgomp.fortran/metadirective-6.f90: New test. + 2023-07-26 Tobias Burnus <tobias@codesourcery.com> Backported from master: diff --git a/libgomp/testsuite/libgomp.fortran/metadirective-1.f90 b/libgomp/testsuite/libgomp.fortran/metadirective-1.f90 index 9f6a07459e0..7b3e09f7c2a 100644 --- a/libgomp/testsuite/libgomp.fortran/metadirective-1.f90 +++ b/libgomp/testsuite/libgomp.fortran/metadirective-1.f90 @@ -17,17 +17,45 @@ program test do i = 1, N if (z(i) .ne. x(i) * y(i)) stop 1 end do + + ! ----- + do i = 1, N + x(i) = i; + y(i) = -i; + end do + + call g (x, y, z) + + do i = 1, N + if (z(i) .ne. x(i) * y(i)) stop 1 + end do + contains subroutine f (x, y, z) integer :: x(N), y(N), z(N) !$omp target map (to: x, y) map(from: z) + block + !$omp metadirective & + !$omp& when(device={arch("nvptx")}: teams loop) & + !$omp& default(parallel loop) + do i = 1, N + z(i) = x(i) * y(i) + enddo + end block + end subroutine + subroutine g (x, y, z) + integer :: x(N), y(N), z(N) + + !$omp target map (to: x, y) map(from: z) + block !$omp metadirective & !$omp& when(device={arch("nvptx")}: teams loop) & !$omp& default(parallel loop) do i = 1, N z(i) = x(i) * y(i) enddo + end block !$omp end target end subroutine end program diff --git a/libgomp/testsuite/libgomp.fortran/metadirective-6.f90 b/libgomp/testsuite/libgomp.fortran/metadirective-6.f90 new file mode 100644 index 00000000000..436fdbade2f --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/metadirective-6.f90 @@ -0,0 +1,58 @@ +! { dg-do compile } + +program test + implicit none + + integer, parameter :: N = 100 + integer :: x(N), y(N), z(N) + integer :: i + +contains + subroutine f (x, y, z) + integer :: x(N), y(N), z(N) + + !$omp target map (to: x, y) map(from: z) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS at .2. may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + block + !$omp metadirective & + !$omp& when(device={arch("nvptx")}: teams loop) & + !$omp& default(parallel loop) ! { dg-error "\\(1\\)" } + ! FIXME: The line above should be the same error as above but some fails here with -fno-diagnostics-show-caret + ! Seems as if some gcc/testsuite/ fix is missing for libgomp/testsuite + do i = 1, N + z(i) = x(i) * y(i) + enddo + z(N) = z(N) + 1 ! <<< invalid + end block + end subroutine + + subroutine f2 (x, y, z) + integer :: x(N), y(N), z(N) + + !$omp target map (to: x, y) map(from: z) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + block + integer :: i ! << invalid + !$omp metadirective & + !$omp& when(device={arch("nvptx")}: teams loop) & + !$omp& default(parallel loop) + do i = 1, N + z(i) = x(i) * y(i) + enddo + end block + end subroutine + subroutine g (x, y, z) + integer :: x(N), y(N), z(N) + + !$omp target map (to: x, y) map(from: z) ! { dg-error "OMP TARGET region at .1. with a nested TEAMS may not contain any other statement, declaration or directive outside of the single TEAMS construct" } + block + !$omp metadirective & ! <<<< invalid + !$omp& when(device={arch("nvptx")}: flush) & + !$omp& default(nothing) + !$omp teams loop + do i = 1, N + z(i) = x(i) * y(i) + enddo + end block + !$omp end target + end subroutine + +end program ^ permalink raw reply [flat|nested] 6+ messages in thread
end of thread, other threads:[~2023-07-27 16:37 UTC | newest] Thread overview: 6+ messages (download: mbox.gz / follow: Atom feed) -- links below jump to the message on this page -- 2023-07-24 19:43 [Patch] OpenMP/Fortran: Reject not strictly nested target -> teams [PR110725, PR71065] Tobias Burnus 2023-07-24 19:49 ` Jakub Jelinek 2023-07-24 20:05 ` Tobias Burnus 2023-07-25 7:37 ` Tobias Burnus 2023-07-25 11:14 ` [patch] OpenMP/Fortran: Reject declarations between target + teams (was: [Patch] OpenMP/Fortran: Reject not strictly nested target -> teams [PR110725, PR71065]) Tobias Burnus 2023-07-27 16:36 ` [committed] OpenMP/Fortran: Extend reject code between target + teams [PR71065, PR110725] (was: Re: [patch] OpenMP/Fortran: Reject declarations between target + teams (was: [Patch] OpenMP/Fortran: Reject not strictly nested target -> teams [PR110725, PR71065])) Tobias Burnus
This is a public inbox, see mirroring instructions for how to clone and mirror all data and code used for this inbox; as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).