Hi! The OpenMP standard says: "A teams region can only be strictly nested within the implicit parallel region or a target region. If a teams construct is nested within a target construct, that target construct must contain no statements, declarations or directives outside of the teams construct." We weren't diagnosing that restriction, because we need to allow e.g. #pragma omp target {{{{{{ #pragma omp teams ; }}}}}} and as target doesn't need to have teams nested in it, using some special parser of the target body didn't feel right. And after the parsing, the question is if e.g. already parsing of the clauses doesn't add some statements before the teams statement (gimplification certainly will).
As we now have a bugreport where we ICE on the invalid code, this just diagnoses a subset of the invalid programs, in particular those where nest to the teams strictly nested in targets the target region contains some other OpenMP construct. Bootstrapped/regtested on x86_64-linux and i686-linux, committed to trunk. 2021-02-24 Jakub Jelinek <ja...@redhat.com> PR fortran/99226 * omp-low.c (struct omp_context): Add teams_nested_p and nonteams_nested_p members. (scan_omp_target): Diagnose teams nested inside of target with other directives strictly nested inside of the same target. (check_omp_nesting_restrictions): Set ctx->teams_nested_p or ctx->nonteams_nested_p as needed. * c-c++-common/gomp/pr99226.c: New test. * gfortran.dg/gomp/pr99226.f90: New test. --- gcc/omp-low.c.jj 2021-01-19 13:26:07.282118987 +0100 +++ gcc/omp-low.c 2021-02-24 15:06:41.897529041 +0100 @@ -171,6 +171,14 @@ struct omp_context /* True if there is bind clause on the construct (i.e. a loop construct). */ bool loop_p; + + /* Only used for omp target contexts. True if a teams construct is + strictly nested in it. */ + bool teams_nested_p; + + /* Only used for omp target contexts. True if an OpenMP construct other + than teams is strictly nested in it. */ + bool nonteams_nested_p; }; static splay_tree all_contexts; @@ -2956,6 +2964,14 @@ scan_omp_target (gomp_target *stmt, omp_ if (offloaded) fixup_child_record_type (ctx); } + + if (ctx->teams_nested_p && ctx->nonteams_nested_p) + { + error_at (gimple_location (stmt), + "%<target%> construct with nested %<teams%> construct " + "contains directives outside of the %<teams%> construct"); + gimple_omp_set_body (stmt, gimple_build_bind (NULL, NULL, NULL)); + } } /* Scan an OpenMP teams directive. */ @@ -3025,6 +3041,14 @@ check_omp_nesting_restrictions (gimple * if (ctx != NULL) { + if (gimple_code (ctx->stmt) == GIMPLE_OMP_TARGET + && gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_REGION) + { + if (gimple_code (stmt) == GIMPLE_OMP_TEAMS && !ctx->teams_nested_p) + ctx->teams_nested_p = true; + else + ctx->nonteams_nested_p = true; + } if (gimple_code (ctx->stmt) == GIMPLE_OMP_SCAN && ctx->outer && gimple_code (ctx->outer->stmt) == GIMPLE_OMP_FOR) --- gcc/testsuite/c-c++-common/gomp/pr99226.c.jj 2021-02-24 15:04:39.213747973 +0100 +++ gcc/testsuite/c-c++-common/gomp/pr99226.c 2021-02-24 15:02:58.068825607 +0100 @@ -0,0 +1,17 @@ +/* PR fortran/99226 */ +/* { dg-do compile } */ + +void +foo (int n) +{ + int i; + #pragma omp target /* { dg-error "construct with nested 'teams' construct contains directives outside of the 'teams' construct" } */ + { + #pragma omp teams distribute dist_schedule (static, n + 4) + for (i = 0; i < 8; i++) + ; + #pragma omp teams distribute dist_schedule (static, n + 4) + for (i = 0; i < 8; i++) + ; + } +} --- gcc/testsuite/gfortran.dg/gomp/pr99226.f90.jj 2021-02-24 15:04:21.994919040 +0100 +++ gcc/testsuite/gfortran.dg/gomp/pr99226.f90 2021-02-24 15:03:55.618185552 +0100 @@ -0,0 +1,13 @@ +! PR fortran/99226 + +subroutine sub (n) + integer :: n, i + !$omp target ! { dg-error "construct with nested 'teams' construct contains directives outside of the 'teams' construct" } + !$omp teams distribute dist_schedule (static,n+4) + do i = 1, 8 + end do + !$omp teams distribute dist_schedule (static,n+4) + do i = 1, 8 + end do + !$omp end target +end Jakub