Hi! For the tree-nested.c purposes, host teams construct is like parallel or task, we need shared clause to propagate the frame structures etc., after all, host teams is just another parallelization level, just spread across NUMA nodes. For the target teams that is not the case, we ignore there all shared clauses because we handle it very specially, it must be immediately nested in target anyway, so the target just starts all the teams.
Bootstrapped/regtested on x86_64-linux and i686-linux, committed to trunk. 2019-01-24 Jakub Jelinek <ja...@redhat.com> PR middle-end/89015 * tree-nested.c (convert_nonlocal_reference_stmt, convert_local_reference_stmt, convert_tramp_reference_stmt, convert_gimple_call) <case GIMPLE_OMP_TEAMS>: Treat gimple_omp_teams_host teams stmts like GIMPLE_OMP_PARALLEL or GIMPLE_OMP_TASK. * gcc.dg/gomp/pr89015.c: New test. --- gcc/tree-nested.c.jj 2019-01-01 12:37:19.771932062 +0100 +++ gcc/tree-nested.c 2019-01-24 17:51:39.817047968 +0100 @@ -1497,6 +1497,20 @@ convert_nonlocal_reference_stmt (gimple_ } break; + case GIMPLE_OMP_TEAMS: + if (!gimple_omp_teams_host (as_a <gomp_teams *> (stmt))) + { + save_suppress = info->suppress_expansion; + convert_nonlocal_omp_clauses (gimple_omp_teams_clauses_ptr (stmt), + wi); + walk_body (convert_nonlocal_reference_stmt, + convert_nonlocal_reference_op, info, + gimple_omp_body_ptr (stmt)); + info->suppress_expansion = save_suppress; + break; + } + /* FALLTHRU */ + case GIMPLE_OMP_PARALLEL: case GIMPLE_OMP_TASK: save_suppress = info->suppress_expansion; @@ -1601,14 +1615,6 @@ convert_nonlocal_reference_stmt (gimple_ info->suppress_expansion = save_suppress; break; - case GIMPLE_OMP_TEAMS: - save_suppress = info->suppress_expansion; - convert_nonlocal_omp_clauses (gimple_omp_teams_clauses_ptr (stmt), wi); - walk_body (convert_nonlocal_reference_stmt, convert_nonlocal_reference_op, - info, gimple_omp_body_ptr (stmt)); - info->suppress_expansion = save_suppress; - break; - case GIMPLE_OMP_SECTION: case GIMPLE_OMP_MASTER: case GIMPLE_OMP_ORDERED: @@ -2168,6 +2174,18 @@ convert_local_reference_stmt (gimple_stm switch (gimple_code (stmt)) { + case GIMPLE_OMP_TEAMS: + if (!gimple_omp_teams_host (as_a <gomp_teams *> (stmt))) + { + save_suppress = info->suppress_expansion; + convert_local_omp_clauses (gimple_omp_teams_clauses_ptr (stmt), wi); + walk_body (convert_local_reference_stmt, convert_local_reference_op, + info, gimple_omp_body_ptr (stmt)); + info->suppress_expansion = save_suppress; + break; + } + /* FALLTHRU */ + case GIMPLE_OMP_PARALLEL: case GIMPLE_OMP_TASK: save_suppress = info->suppress_expansion; @@ -2299,14 +2317,6 @@ convert_local_reference_stmt (gimple_stm info->static_chain_added |= save_static_chain_added; break; - case GIMPLE_OMP_TEAMS: - save_suppress = info->suppress_expansion; - convert_local_omp_clauses (gimple_omp_teams_clauses_ptr (stmt), wi); - walk_body (convert_local_reference_stmt, convert_local_reference_op, - info, gimple_omp_body_ptr (stmt)); - info->suppress_expansion = save_suppress; - break; - case GIMPLE_OMP_SECTION: case GIMPLE_OMP_MASTER: case GIMPLE_OMP_ORDERED: @@ -2607,6 +2617,14 @@ convert_tramp_reference_stmt (gimple_stm break; } + case GIMPLE_OMP_TEAMS: + if (!gimple_omp_teams_host (as_a <gomp_teams *> (stmt))) + { + *handled_ops_p = false; + return NULL_TREE; + } + goto do_parallel; + case GIMPLE_OMP_TARGET: if (!is_gimple_omp_offloaded (stmt)) { @@ -2616,6 +2634,7 @@ convert_tramp_reference_stmt (gimple_stm /* FALLTHRU */ case GIMPLE_OMP_PARALLEL: case GIMPLE_OMP_TASK: + do_parallel: { tree save_local_var_chain = info->new_local_var_chain; walk_gimple_op (stmt, convert_tramp_reference_op, wi); @@ -2723,6 +2742,15 @@ convert_gimple_call (gimple_stmt_iterato } break; + case GIMPLE_OMP_TEAMS: + if (!gimple_omp_teams_host (as_a <gomp_teams *> (stmt))) + { + walk_body (convert_gimple_call, NULL, info, + gimple_omp_body_ptr (stmt)); + break; + } + /* FALLTHRU */ + case GIMPLE_OMP_PARALLEL: case GIMPLE_OMP_TASK: save_static_chain_added = info->static_chain_added; @@ -2798,7 +2826,6 @@ convert_gimple_call (gimple_stmt_iterato case GIMPLE_OMP_SECTIONS: case GIMPLE_OMP_SECTION: case GIMPLE_OMP_SINGLE: - case GIMPLE_OMP_TEAMS: case GIMPLE_OMP_MASTER: case GIMPLE_OMP_TASKGROUP: case GIMPLE_OMP_ORDERED: --- gcc/testsuite/gcc.dg/gomp/pr89015.c.jj 2019-01-24 17:56:18.018471801 +0100 +++ gcc/testsuite/gcc.dg/gomp/pr89015.c 2019-01-24 17:56:02.735722800 +0100 @@ -0,0 +1,25 @@ +/* PR middle-end/89015 */ +/* { dg-do compile } */ + +int +foo (int n, float *x, float *y) +{ + int i; + int bar (void) { return i; } +#pragma omp teams distribute parallel for simd + for (i = 0; i < n; i++) + y[i] = x[i]; + return bar (); +} + +int +baz (int n, float *x, float *y) +{ + int i; + int qux (void) { +#pragma omp teams distribute parallel for simd + for (i = 0; i < n; i++) + y[i] = x[i]; + } + return qux (); +} Jakub