Short version:
* dyn_groupprivate clause - valid on target and on teams
Fixed: was not accepted on 'teams'
Fixed: on combined directives, the clause was lost.
* target + teams (two directives or in one combined directive):
For the dyn_groupprivate, num_teams, thread_limit clauses:
- now evaluated on the host
- clauses are put on 'target' (also num_teams, which only
accepts this clause internally)
- for dyn_groupprivate/thread_limit, the value on 'teams' wins
if both target and teams such a clause
- there will be no dyn_groupprivate on teams
- num_teams + thread_limit: temp variable will be firstprivate
on 'target' and clauses are also on 'teams'
Background: for all three clauses, the host code launching a
GPU kernel should know the value at launch time. If not, it
the value has to be guesses. - OpenMP 5.0 and then 6.0 added
wording to make it possible to evaluate it on the host, in time
for the kernel launch.
(Currently, num_teams/num_threads == -1 is used in the call to
GOMP_target_ext to indicate that a teams with a clause exists
but the value is unknown when GOMP_target_ext is called.
With this patch, gfortran code shall never yield '-1' any more.)
Any comments before I commit this later today?
Tobias
* * *
Long version for the second change:
When launching a GPU kernel, the host code needs to know how many
teams and threads have to be created. That's not that trivial for
code like:
!$omp target
!$omp teams num_teams(f()) thread_limit(n[0])
as f() on the host or on the device might return different values
and 'n' might be mapped on the target - such that evaluation on
host and device can be again have different values.
Since 5.0, OpenMP has clarified that on 'target teams' the evaluation
is always on the host - and, since 6.0, also for code like above,
OpenMP requires the user that evaluation on the host and on the device
will yield the same result (target-consistent expression).
[No sane person will code with non-target-consistent expressions.]
There is code in the middle end which is supposed to handle code as
above in simple cases like:
integer :: n
n = 16
!$omp target
!$omp teams num_teams(n)
or even
!$omp target teams num_teams(n)
However, before this patch is turned into:
#pragma omp target num_teams(-1) firstprivate(n)
{
int D.123 = n; // expanded by the front end
#pragma teams num_teams(D.123)
as the middle end expects the full expression in the clause
(as common with C) and fails due to the temporary.
Now the result with this patch is the front-end code:
int D.123 = n;
#pragma omp target num_teams(D.123) firstprivate(D.123)
{
#pragma teams num_teams(D.123)
and the middle change of this patch ensures that this
remains such after gimplification without adding a
second num_teams(-1) to 'target'.
Tobias
PS: For C/C++, the situation is better; for the combined directive,
the host expansion and firstprivate already happens (5.0 semantic)
and the gimplify.cc code will also help with the common cases when
using two directives. - Still, a similar change for C/C++ should
be done in lockstep with supporting spatial dimensions to completely
avoid the '-1' case. (target teams without clauses is '0' and target
without teams is '-2'.)
OpenMP/Fortran: Fixes to target-consistent clauses and dyn_groupprivate
num_teams and thread_limit clauses on teams inside target can be evaluated
on the host (and must be the same as on the device: target consistent).
As those values are relevant when launching a GPU kernel, it very much
makes sense to do so. This commit ensures that this happens (gimplify.cc's
optimize_target_teams comes too late as the expanded expression has been
converted to assignment statements - likewise for combined directives
and for immediately nested directives.)
Note: Since 5.0, the evaluation of the num_teams and thread_limit on a
combined directive of target and teams must be evaluated on the host,
since 6.0 this is expanded to also cover a teams construct that is
immediately nested in target. - This is is a breaking change as before
a function call to a device function instead to a host function or a
mapped variable instead of the corresponding host function could yield
other results when evaluated on the host or on the device. However,
this is extremely unlikely to occur in real-world code.
This commit evaluates the expression of those two clauses and of
dyn_groupprivate on the host; dyn_groupprivate will only be put on
target, the other clauses are on both target and teams and the
temporary clause will also be put into a firstprivate clause on target.
(num_teams is not permitted on target by the standard but internally,
GCC puts it there, too.)
For dyn_groupprivate and thread_limit, the clause on teams wins
as it overrides the settings of target - such that it can be used directly.
NOTE: 'if' and 'allocate' are also target consistent; for 'allocate',
this is no restriction due to other restrictions - and 'if' on teams
is not yet supported.
gcc/fortran/ChangeLog:
* openmp.cc (OMP_TEAMS_CLAUSES): Add OMP_CLAUSE_DYN_GROUPPRIVATE.
* trans-openmp.cc (gfc_trans_omp_clauses): Use expression and not
directive line for the location.
(gfc_split_omp_clauses): Handle dyn_groupprivate.
(gfc_trans_omp_target): Handle target-consistent clauses.
gcc/ChangeLog:
* gimplify.cc (optimize_target_teams): SKip if there is already a
'num_teams' (i.e. already optimized by the front end).
gcc/testsuite/ChangeLog:
* gfortran.dg/gomp/spatial-dimensions-1.f90: Update scan-tree-dump.
* gfortran.dg/gomp/spatial-dimensions-4.f90: Likewise.
* gfortran.dg/gomp/spatial-dimensions-5.f90: Likewise.
* gfortran.dg/gomp/dyn_groupprivate-3.f90: New test.
* gfortran.dg/gomp/dyn_groupprivate-4.f90: New test.
* gfortran.dg/gomp/target-teams-1.f90: New test.
* gfortran.dg/gomp/target-teams-2.f90: New test.
gcc/fortran/openmp.cc | 2 +-
gcc/fortran/trans-openmp.cc | 233 +++++++++++++++++----
gcc/gimplify.cc | 6 +
.../gfortran.dg/gomp/dyn_groupprivate-3.f90 | 24 +++
.../gfortran.dg/gomp/dyn_groupprivate-4.f90 | 184 ++++++++++++++++
.../gfortran.dg/gomp/spatial-dimensions-1.f90 | 2 +-
.../gfortran.dg/gomp/spatial-dimensions-4.f90 | 15 +-
.../gfortran.dg/gomp/spatial-dimensions-5.f90 | 12 +-
gcc/testsuite/gfortran.dg/gomp/target-teams-1.f90 | 65 ++++++
gcc/testsuite/gfortran.dg/gomp/target-teams-2.f90 | 17 ++
10 files changed, 504 insertions(+), 56 deletions(-)
diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index ba7f3c5ea24..670e4af1399 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -5995,7 +5995,7 @@ cleanup:
(omp_mask (OMP_CLAUSE_NUM_TEAMS) | OMP_CLAUSE_THREAD_LIMIT \
| OMP_CLAUSE_DEFAULT | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE \
| OMP_CLAUSE_SHARED | OMP_CLAUSE_REDUCTION | OMP_CLAUSE_ALLOCATE \
- | OMP_CLAUSE_MESSAGE)
+ | OMP_CLAUSE_MESSAGE | OMP_CLAUSE_DYN_GROUPPRIVATE)
#define OMP_DISTRIBUTE_CLAUSES \
(omp_mask (OMP_CLAUSE_PRIVATE) | OMP_CLAUSE_FIRSTPRIVATE \
| OMP_CLAUSE_LASTPRIVATE | OMP_CLAUSE_COLLAPSE | OMP_CLAUSE_DIST_SCHEDULE \
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index bee9fd608ec..e64aba9aef7 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -5760,8 +5760,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
case OMP_FALLBACK_NONE:
break;
}
- c = build_omp_clause (gfc_get_location (&where),
- OMP_CLAUSE_DYN_GROUPPRIVATE);
+ location_t loc = gfc_get_location (&clauses->dyn_groupprivate->where);
+ c = build_omp_clause (loc, OMP_CLAUSE_DYN_GROUPPRIVATE);
OMP_CLAUSE_DYN_GROUPPRIVATE_KIND (c) = kind;
OMP_CLAUSE_DYN_GROUPPRIVATE_EXPR (c) = expr;
omp_clauses = gfc_trans_add_clause (c, omp_clauses);
@@ -6016,9 +6016,19 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
gfc_convert_expr_to_tree (block, el->expr),
num_teams);
c = build_omp_clause (gfc_get_location (&where), OMP_CLAUSE_NUM_TEAMS);
- OMP_CLAUSE_NUM_TEAMS_LOWER_EXPR (c) = NULL_TREE;
- OMP_CLAUSE_NUM_TEAMS_UPPER_EXPR (c) = nreverse (num_teams);
- OMP_CLAUSE_NUM_TEAMS_DIMS (c) = clauses->num_teams_dims;
+ if (clauses->num_teams_dims)
+ {
+ OMP_CLAUSE_NUM_TEAMS_LOWER_EXPR (c) = NULL_TREE;
+ OMP_CLAUSE_NUM_TEAMS_UPPER_EXPR (c) = nreverse (num_teams);
+ OMP_CLAUSE_NUM_TEAMS_DIMS (c) = clauses->num_teams_dims;
+ }
+ else
+ {
+ OMP_CLAUSE_NUM_TEAMS_UPPER_EXPR (c) = TREE_VALUE (num_teams);
+ if (TREE_CHAIN(num_teams))
+ OMP_CLAUSE_NUM_TEAMS_LOWER_EXPR (c)
+ = TREE_VALUE (TREE_CHAIN (num_teams));
+ }
omp_clauses = gfc_trans_add_clause (c, omp_clauses);
}
@@ -8565,6 +8575,10 @@ gfc_split_omp_clauses (gfc_code *code,
= code->ext.omp_clauses->lists[OMP_LIST_IS_DEVICE_PTR];
clausesa[GFC_OMP_SPLIT_TARGET].lists[OMP_LIST_HAS_DEVICE_ADDR]
= code->ext.omp_clauses->lists[OMP_LIST_HAS_DEVICE_ADDR];
+ clausesa[GFC_OMP_SPLIT_TARGET].dyn_groupprivate
+ = code->ext.omp_clauses->dyn_groupprivate;
+ clausesa[GFC_OMP_SPLIT_TARGET].fallback
+ = code->ext.omp_clauses->fallback;
clausesa[GFC_OMP_SPLIT_TARGET].device
= code->ext.omp_clauses->device;
clausesa[GFC_OMP_SPLIT_TARGET].thread_limit_list
@@ -8603,6 +8617,10 @@ gfc_split_omp_clauses (gfc_code *code,
= code->ext.omp_clauses->thread_limit_strict;
clausesa[GFC_OMP_SPLIT_TEAMS].thread_limit_dims
= code->ext.omp_clauses->thread_limit_dims;
+ clausesa[GFC_OMP_SPLIT_TEAMS].dyn_groupprivate
+ = code->ext.omp_clauses->dyn_groupprivate;
+ clausesa[GFC_OMP_SPLIT_TEAMS].fallback
+ = code->ext.omp_clauses->fallback;
/* Shared and default clauses are allowed on parallel, teams
and taskloop. */
clausesa[GFC_OMP_SPLIT_TEAMS].lists[OMP_LIST_SHARED]
@@ -9405,16 +9423,181 @@ gfc_trans_omp_target (gfc_code *code)
tree stmt, omp_clauses = NULL_TREE;
gfc_start_block (&block);
+
+ /* For target + teams, evaluate some clauses on the host; as for those their
+ expressions must be target consistent, this is also permitted when
+ nested and not only for combined directives. */
+ bool combined_teams_dir = false;
+ if (flag_openmp)
+ switch (code->op)
+ {
+ case EXEC_OMP_TARGET_TEAMS:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+ case EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD:
+ case EXEC_OMP_TARGET_TEAMS_LOOP:
+ combined_teams_dir = true;
+ break;
+ default:
+ combined_teams_dir = false;
+ break;
+ }
+ tree tree_tgt_clauses = NULL_TREE, tree_teams_clauses = NULL_TREE;
+ gfc_omp_clauses *teams_clauses = NULL;
+ if (flag_openmp && code->op == EXEC_OMP_TARGET && code->block->next)
+ {
+ gfc_code *code2 = (code->block->next->op == EXEC_BLOCK
+ ? code->block->next->ext.block.ns->code
+ : code->block->next);
+ if (code2)
+ switch (code2->op)
+ {
+ case EXEC_OMP_TEAMS:
+ case EXEC_OMP_TEAMS_DISTRIBUTE:
+ case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO:
+ case EXEC_OMP_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD:
+ case EXEC_OMP_TEAMS_DISTRIBUTE_SIMD:
+ case EXEC_OMP_TEAMS_LOOP:
+ teams_clauses = code2->ext.omp_clauses;
+ break;
+ default:
+ break;
+ }
+ }
+ if ((teams_clauses && teams_clauses->num_teams_list)
+ || (combined_teams_dir && code->ext.omp_clauses->num_teams_list))
+ {
+ gfc_omp_clauses *num_teams_clauses
+ = teams_clauses ? teams_clauses : code->ext.omp_clauses;
+ tree c, t, num_teams = NULL_TREE;
+ for (gfc_expr_list *el = num_teams_clauses->num_teams_list;
+ el; el = el->next)
+ {
+ t = gfc_convert_expr_to_tree (&block, el->expr);
+ if (DECL_P (t))
+ {
+ c = build_omp_clause (gfc_get_location (&el->expr->where),
+ OMP_CLAUSE_FIRSTPRIVATE);
+ OMP_CLAUSE_DECL (c) = t;
+ tree_tgt_clauses = gfc_trans_add_clause (c, tree_tgt_clauses);
+ }
+ num_teams = tree_cons (NULL_TREE, t, num_teams);
+ }
+ locus *loc = &num_teams_clauses->num_teams_list->expr->where;
+ c = build_omp_clause (gfc_get_location (loc), OMP_CLAUSE_NUM_TEAMS);
+ if (num_teams_clauses->num_teams_dims)
+ {
+ OMP_CLAUSE_NUM_TEAMS_LOWER_EXPR (c) = NULL_TREE;
+ OMP_CLAUSE_NUM_TEAMS_UPPER_EXPR (c) = nreverse (num_teams);
+ OMP_CLAUSE_NUM_TEAMS_DIMS (c) = num_teams_clauses->num_teams_dims;
+ }
+ else
+ {
+ OMP_CLAUSE_NUM_TEAMS_UPPER_EXPR (c) = TREE_VALUE (num_teams);
+ if (TREE_CHAIN (num_teams))
+ OMP_CLAUSE_NUM_TEAMS_LOWER_EXPR (c)
+ = TREE_VALUE (TREE_CHAIN (num_teams));
+ }
+ tree_tgt_clauses = gfc_trans_add_clause (c, tree_tgt_clauses);
+ tree_teams_clauses = gfc_trans_add_clause (unshare_expr (c),
+ tree_teams_clauses);
+ gfc_free_expr_list (num_teams_clauses->num_teams_list);
+ num_teams_clauses->num_teams_list = NULL;
+ }
+ if ((teams_clauses && teams_clauses->thread_limit_list)
+ || (combined_teams_dir && code->ext.omp_clauses->thread_limit_list))
+ {
+ /* thread_limit can be put on target and on teams, use the team one as
+ having higher precedence. */
+ gfc_omp_clauses *thread_limit_clauses
+ = (teams_clauses && teams_clauses->thread_limit_list
+ ? teams_clauses : code->ext.omp_clauses);
+ tree c, t, thread_limit = NULL_TREE;
+ for (gfc_expr_list *el = thread_limit_clauses->thread_limit_list;
+ el; el = el->next)
+ {
+ t = gfc_convert_expr_to_tree (&block, el->expr);
+ if (DECL_P (t))
+ {
+ c = build_omp_clause (gfc_get_location (&el->expr->where),
+ OMP_CLAUSE_FIRSTPRIVATE);
+ OMP_CLAUSE_DECL (c) = t;
+ tree_tgt_clauses = gfc_trans_add_clause (c, tree_tgt_clauses);
+ }
+ thread_limit = tree_cons (NULL_TREE, t, thread_limit);
+ }
+ locus *loc = &thread_limit_clauses->thread_limit_list->expr->where;
+ c = build_omp_clause (gfc_get_location (loc), OMP_CLAUSE_THREAD_LIMIT);
+ OMP_CLAUSE_THREAD_LIMIT_EXPR (c) = nreverse (thread_limit);
+ OMP_CLAUSE_THREAD_LIMIT_STRICT (c)
+ = thread_limit_clauses->thread_limit_strict;
+ OMP_CLAUSE_THREAD_LIMIT_DIMS (c)
+ = thread_limit_clauses->thread_limit_dims;
+ tree_tgt_clauses = gfc_trans_add_clause (c, tree_tgt_clauses);
+ tree_teams_clauses = gfc_trans_add_clause (unshare_expr (c),
+ tree_teams_clauses);
+ gfc_free_expr_list (thread_limit_clauses->thread_limit_list);
+ thread_limit_clauses->thread_limit_list = NULL;
+ gfc_free_expr_list (code->ext.omp_clauses->thread_limit_list);
+ code->ext.omp_clauses->thread_limit_list = NULL;
+ }
+ if ((teams_clauses && teams_clauses->dyn_groupprivate)
+ || (combined_teams_dir && code->ext.omp_clauses->dyn_groupprivate))
+ {
+ /* dyn_groupprivate can be put on target and on teams, use the team one as
+ having higher precedence. */
+ gfc_omp_clauses *dyn_gp_clauses
+ = (teams_clauses && teams_clauses->dyn_groupprivate
+ ? teams_clauses : code->ext.omp_clauses);
+ gfc_omp_clauses tmp = {};
+ tmp.dyn_groupprivate = dyn_gp_clauses->dyn_groupprivate;
+ tmp.fallback = dyn_gp_clauses->fallback;
+ tree c = gfc_trans_omp_clauses (&block, &tmp, code->loc);
+ if (DECL_P (OMP_CLAUSE_DYN_GROUPPRIVATE_EXPR (c)))
+ {
+ tree c2;
+ OMP_CLAUSE_DYN_GROUPPRIVATE_EXPR (c)
+ = gfc_evaluate_now (OMP_CLAUSE_DYN_GROUPPRIVATE_EXPR (c), &block);
+ c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+ OMP_CLAUSE_FIRSTPRIVATE);
+ OMP_CLAUSE_DECL (c2) = OMP_CLAUSE_DYN_GROUPPRIVATE_EXPR (c);
+ tree_tgt_clauses = gfc_trans_add_clause (c2, tree_tgt_clauses);
+ }
+ /* There is no need to put the clause also on teams. */
+ tree_tgt_clauses = gfc_trans_add_clause (c, tree_tgt_clauses);
+ gfc_free_expr (dyn_gp_clauses->dyn_groupprivate);
+ dyn_gp_clauses->dyn_groupprivate = NULL;
+ gfc_free_expr (code->ext.omp_clauses->dyn_groupprivate);
+ code->ext.omp_clauses->dyn_groupprivate = NULL;
+ }
+ /* FIXME: also 'if' and 'allocate' are target consistent. */
+
gfc_split_omp_clauses (code, clausesa);
if (flag_openmp)
omp_clauses
- = gfc_trans_omp_clauses (&block, &clausesa[GFC_OMP_SPLIT_TARGET],
- code->loc);
+ = chainon (gfc_trans_omp_clauses (&block, &clausesa[GFC_OMP_SPLIT_TARGET],
+ code->loc), tree_tgt_clauses);
switch (code->op)
{
case EXEC_OMP_TARGET:
pushlevel ();
- stmt = gfc_trans_omp_code (code->block->next, true);
+ if (tree_teams_clauses)
+ {
+ gfc_code *code2 = code->block->next;
+ if (code2->op == EXEC_BLOCK)
+ {
+ gcc_checking_assert (!code2->ext.block.ns->sym_root->left
+ && !code2->ext.block.ns->sym_root->right);
+ code2 = code2->ext.block.ns->code;
+ }
+ gcc_checking_assert (code2->next == NULL);
+ stmt = gfc_trans_omp_teams (code2,
+ combined_teams_dir ? clausesa : NULL,
+ tree_teams_clauses);
+ }
+ else
+ stmt = gfc_trans_omp_code (code->block->next, true);
stmt = build3_v (BIND_EXPR, NULL, stmt, poplevel (1, 0));
break;
case EXEC_OMP_TARGET_PARALLEL:
@@ -9464,38 +9647,8 @@ gfc_trans_omp_target (gfc_code *code)
poplevel (0, 0);
break;
default:
- if (flag_openmp
- && (clausesa[GFC_OMP_SPLIT_TEAMS].num_teams_list
- || clausesa[GFC_OMP_SPLIT_TEAMS].thread_limit_list))
- {
- gfc_omp_clauses clausesb;
- tree teams_clauses;
- /* For combined !$omp target teams, the num_teams and
- thread_limit clauses are evaluated before entering the
- target construct. */
- memset (&clausesb, '\0', sizeof (clausesb));
- clausesb.num_teams_list
- = clausesa[GFC_OMP_SPLIT_TEAMS].num_teams_list;
- clausesb.num_teams_dims
- = clausesa[GFC_OMP_SPLIT_TEAMS].num_teams_dims;
- clausesb.thread_limit_list
- = clausesa[GFC_OMP_SPLIT_TEAMS].thread_limit_list;
- clausesb.thread_limit_strict
- = clausesa[GFC_OMP_SPLIT_TEAMS].thread_limit_strict;
- clausesb.thread_limit_dims
- = clausesa[GFC_OMP_SPLIT_TEAMS].thread_limit_dims;
- clausesa[GFC_OMP_SPLIT_TEAMS].num_teams_list = NULL;
- clausesa[GFC_OMP_SPLIT_TEAMS].thread_limit_list = NULL;
- teams_clauses
- = gfc_trans_omp_clauses (&block, &clausesb, code->loc);
- pushlevel ();
- stmt = gfc_trans_omp_teams (code, clausesa, teams_clauses);
- }
- else
- {
- pushlevel ();
- stmt = gfc_trans_omp_teams (code, clausesa, NULL_TREE);
- }
+ pushlevel ();
+ stmt = gfc_trans_omp_teams (code, clausesa, tree_teams_clauses);
if (TREE_CODE (stmt) != BIND_EXPR)
stmt = build3_v (BIND_EXPR, NULL, stmt, poplevel (1, 0));
else
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 94d71d71fc2..d521bda8186 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -18790,6 +18790,12 @@ computable_teams_clause (tree *tp, int *walk_subtrees, void *)
static void
optimize_target_teams (tree target, gimple_seq *pre_p)
{
+ /* If the NUM_TEAMS clause is on TARGET, assume that the front end already
+ took care of the target-consistent clauses. Note that user code cannot
+ specify 'num_teams' on 'target', only on 'teams'. */
+ if (omp_find_clause (OMP_TARGET_CLAUSES (target), OMP_CLAUSE_NUM_TEAMS))
+ return;
+
tree body = OMP_BODY (target);
tree teams = walk_tree (&body, find_omp_teams, NULL, NULL);
tree num_teams_lower = NULL_TREE;
diff --git a/gcc/testsuite/gfortran.dg/gomp/dyn_groupprivate-3.f90 b/gcc/testsuite/gfortran.dg/gomp/dyn_groupprivate-3.f90
new file mode 100644
index 00000000000..2df14a2c1c9
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/dyn_groupprivate-3.f90
@@ -0,0 +1,24 @@
+! { dg-additional-options "-fdump-tree-original" }
+subroutine foo
+ implicit none
+
+ !$omp teams dyn_groupprivate(fallback(abort) : 123) ! { dg-message "sorry, unimplemented: 'dyn_groupprivate' clause" }
+ !$omp end teams
+
+ !$omp target
+ !$omp teams dyn_groupprivate(fallback(null) : 234) ! { dg-message "sorry, unimplemented: 'dyn_groupprivate' clause" }
+ !$omp end teams
+ !$omp end target
+
+ !$omp teams dyn_groupprivate(fallback(default_mem) : 345) ! { dg-message "sorry, unimplemented: 'dyn_groupprivate' clause" }
+ !$omp end teams
+
+ !$omp target teams dyn_groupprivate(cgroup: 567) ! { dg-message "sorry, unimplemented: 'dyn_groupprivate' clause" }
+ !$omp end target teams
+end
+
+! { dg-final { scan-tree-dump-times "#pragma omp teams dyn_groupprivate\\(fallback\\(abort\\):123\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target dyn_groupprivate\\(fallback\\(null\\):234\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp teams dyn_groupprivate\\(fallback\\(default_mem\\):345\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target dyn_groupprivate\\(567\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp teams\[\\n\\r\]" 2 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/dyn_groupprivate-4.f90 b/gcc/testsuite/gfortran.dg/gomp/dyn_groupprivate-4.f90
new file mode 100644
index 00000000000..0a767c8bec9
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/dyn_groupprivate-4.f90
@@ -0,0 +1,184 @@
+! { dg-additional-options "-fdump-tree-original" }
+
+subroutine f(n)
+ integer, value :: n
+ integer :: i
+
+ ! Clauses only on target (parallel also accepts 'severity')
+
+ !$omp target parallel do simd thread_limit(dims(2), strict : 1,2) dyn_groupprivate(fallback(null), cgroup : 123) message("one") severity(warning) ! { dg-message "sorry, unimplemented: 'dyn_groupprivate' clause" }
+ do i = 1, n
+ end do
+! { dg-final { scan-tree-dump-times "#pragma omp target message\\(\"one\" \\\[len:3\\\]\\) severity\\(fatal\\) dyn_groupprivate\\(fallback\\(null\\):123\\) thread_limit\\(strict,dims\\(\\):1, 2\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp parallel message\\(\"one\" \\\[len:3\\\]\\) severity\\(warning\\)" 1 "original" } }
+
+ !$omp target parallel do thread_limit(dims(2), relaxed : 2,3) dyn_groupprivate(fallback(abort), cgroup : 234) message("two") severity(warning) ! { dg-message "sorry, unimplemented: 'dyn_groupprivate' clause" }
+ do i = 1, n
+ end do
+! { dg-final { scan-tree-dump-times "#pragma omp target message\\(\"two\" \\\[len:3\\\]\\) severity\\(fatal\\) dyn_groupprivate\\(fallback\\(abort\\):234\\) thread_limit\\(dims\\(\\):2, 3\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp parallel message\\(\"two\" \\\[len:3\\\]\\) severity\\(warning\\)" 1 "original" } }
+
+ !$omp target parallel loop thread_limit(relaxed: 34) dyn_groupprivate(fallback(default_mem), cgroup : 345) message("three") severity(warning) ! { dg-message "sorry, unimplemented: 'dyn_groupprivate' clause" }
+ do i = 1, n
+ end do
+! { dg-final { scan-tree-dump-times "#pragma omp target message\\(\"three\" \\\[len:5\\\]\\) severity\\(fatal\\) dyn_groupprivate\\(fallback\\(default_mem\\):345\\) thread_limit\\(34\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp parallel message\\(\"three\" \\\[len:5\\\]\\) severity\\(warning\\)" 1 "original" } }
+
+ !$omp target parallel thread_limit(strict : 45) dyn_groupprivate(cgroup : 456) message("four") severity(warning) ! { dg-message "sorry, unimplemented: 'dyn_groupprivate' clause" }
+ block
+ end block
+! { dg-final { scan-tree-dump-times "#pragma omp target message\\(\"four\" \\\[len:4\\\]\\) severity\\(fatal\\) dyn_groupprivate\\(456\\) thread_limit\\(strict:45\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp parallel message\\(\"four\" \\\[len:4\\\]\\) severity\\(warning\\)" 1 "original" } }
+
+ !$omp target simd thread_limit(dims(1) : 67) dyn_groupprivate(fallback(abort), cgroup : 678) message("five") ! { dg-message "sorry, unimplemented: 'dyn_groupprivate' clause" }
+ do i = 1, n
+ end do
+! { dg-final { scan-tree-dump-times "#pragma omp target message\\(\"five\" \\\[len:4\\\]\\) severity\\(fatal\\) dyn_groupprivate\\(fallback\\(abort\\):678\\) thread_limit\\(strict,dims\\(\\):67\\)" 1 "original" } }
+
+ ! Clauses on target teams (additionally with num_teams)
+
+ !$omp target teams distribute parallel do simd num_teams(11:22) thread_limit(strict : 331) dyn_groupprivate( cgroup ,fallback(null) : 243) message("seven") severity(warning) ! { dg-message "sorry, unimplemented: 'dyn_groupprivate' clause" }
+ do i = 1, n
+ end do
+! { dg-final { scan-tree-dump-times "#pragma omp target message\\(\"seven\" \\\[len:5\\\]\\) severity\\(fatal\\) dyn_groupprivate\\(fallback\\(null\\):243\\) thread_limit\\(strict:331\\) num_teams\\(11:22\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp teams thread_limit\\(strict:331\\) num_teams\\(11:22\\) message\\(\"seven\" \\\[len:5\\\]\\) severity\\(fatal\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp parallel message\\(\"seven\" \\\[len:5\\\]\\) severity\\(warning\\)" 1 "original" } }
+
+ !$omp target teams distribute parallel do num_teams(dims(1):1343) thread_limit(relaxed : 1452) dyn_groupprivate( cgroup ,fallback(default_mem) : 1248) message("nine") severity(warning) ! { dg-message "sorry, unimplemented: 'dyn_groupprivate' clause" }
+ do i = 1, n
+ end do
+! { dg-final { scan-tree-dump-times "#pragma omp target message\\(\"nine\" \\\[len:4\\\]\\) severity\\(fatal\\) dyn_groupprivate\\(fallback\\(default_mem\\):1248\\) thread_limit\\(1452\\) num_teams\\(dims\\(\\):1343\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp teams thread_limit\\(1452\\) num_teams\\(dims\\(\\):1343\\) message\\(\"nine\" \\\[len:4\\\]\\) severity\\(fatal\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp parallel message\\(\"nine\" \\\[len:4\\\]\\) severity\\(warning\\)" 1 "original" } }
+
+ !$omp target teams distribute simd num_teams(dims(2):3,3) thread_limit(938) dyn_groupprivate( 1458) message("ten") ! { dg-message "sorry, unimplemented: 'dyn_groupprivate' clause" }
+ do i = 1, n
+ end do
+! { dg-final { scan-tree-dump-times "#pragma omp target message\\(\"ten\" \\\[len:3\\\]\\) severity\\(fatal\\) dyn_groupprivate\\(1458\\) thread_limit\\(938\\) num_teams\\(dims\\(\\):3, 3\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp teams thread_limit\\(938\\) num_teams\\(dims\\(\\):3, 3\\) message\\(\"ten\" \\\[len:3\\\]\\) severity\\(fatal\\)" 1 "original" } }
+
+ !$omp target teams distribute num_teams(4223) thread_limit(729) dyn_groupprivate( 803) message("eleven") ! { dg-message "sorry, unimplemented: 'dyn_groupprivate' clause" }
+ do i = 1, n
+ end do
+! { dg-final { scan-tree-dump-times "#pragma omp target message\\(\"eleven\" \\\[len:6\\\]\\) severity\\(fatal\\) dyn_groupprivate\\(803\\) thread_limit\\(729\\) num_teams\\(4223\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp teams thread_limit\\(729\\) num_teams\\(4223\\) message\\(\"eleven\" \\\[len:6\\\]\\) severity\\(fatal\\)" 1 "original" } }
+!
+ !$omp target teams loop num_teams(444) thread_limit(451) dyn_groupprivate( 842) message("twelf") ! { dg-message "sorry, unimplemented: 'dyn_groupprivate' clause" }
+ do i = 1, n
+ end do
+! { dg-final { scan-tree-dump-times "#pragma omp target message\\(\"twelf\" \\\[len:5\\\]\\) severity\\(fatal\\) dyn_groupprivate\\(842\\) thread_limit\\(451\\) num_teams\\(444\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp teams thread_limit\\(451\\) num_teams\\(444\\) message\\(\"twelf\" \\\[len:5\\\]\\) severity\\(fatal\\)" 1 "original" } }
+
+
+ ! Same, but in two directives - put on 'teams'
+
+ !$omp target
+ !$omp teams distribute parallel do simd num_teams(11:22) thread_limit(strict : 331) dyn_groupprivate( cgroup ,fallback(null) : 243) message("13") severity(warning) ! { dg-message "sorry, unimplemented: 'dyn_groupprivate' clause" }
+ do i = 1, n
+ end do
+ !$omp end target
+! { dg-final { scan-tree-dump-times "#pragma omp target dyn_groupprivate\\(fallback\\(null\\):243\\) thread_limit\\(strict:331\\) num_teams\\(11:22\\)\[\\n\\r\{ \]+#pragma omp teams thread_limit\\(strict:331\\) num_teams\\(11:22\\) message\\(\"13\" \\\[len:2\\\]\\) severity\\(fatal\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp parallel message\\(\"13\" \\\[len:2\\\]\\) severity\\(warning\\)" 1 "original" } }
+
+ !$omp target
+ !$omp teams distribute parallel do num_teams(dims(1):1343) thread_limit(relaxed : 1452) dyn_groupprivate( cgroup ,fallback(default_mem) : 1248) message("14") severity(warning) ! { dg-message "sorry, unimplemented: 'dyn_groupprivate' clause" }
+ do i = 1, n
+ end do
+ !$omp end target
+! { dg-final { scan-tree-dump-times "#pragma omp target dyn_groupprivate\\(fallback\\(default_mem\\):1248\\) thread_limit\\(1452\\) num_teams\\(dims\\(\\):1343\\)\[\\n\\r\{ \]+#pragma omp teams thread_limit\\(1452\\) num_teams\\(dims\\(\\):1343\\) message\\(\"14\" \\\[len:2\\\]\\) severity\\(fatal\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp parallel message\\(\"14\" \\\[len:2\\\]\\) severity\\(warning\\)" 1 "original" } }
+
+ !$omp target
+ !$omp teams distribute simd num_teams(dims(2):3,3) thread_limit(938) dyn_groupprivate( 1458) message("15") ! { dg-message "sorry, unimplemented: 'dyn_groupprivate' clause" }
+ do i = 1, n
+ end do
+ !$omp end target
+! { dg-final { scan-tree-dump-times "#pragma omp target dyn_groupprivate\\(1458\\) thread_limit\\(938\\) num_teams\\(dims\\(\\):3, 3\\)\[\\n\\r\{ \]+#pragma omp teams thread_limit\\(938\\) num_teams\\(dims\\(\\):3, 3\\) message\\(\"15\" \\\[len:2\\\]\\) severity\\(fatal\\)" 1 "original" } }
+
+ !$omp target
+ !$omp teams distribute num_teams(4223) thread_limit(729) dyn_groupprivate( 803) message("16") ! { dg-message "sorry, unimplemented: 'dyn_groupprivate' clause" }
+ do i = 1, n
+ end do
+ !$omp end target
+! { dg-final { scan-tree-dump-times "#pragma omp target dyn_groupprivate\\(803\\) thread_limit\\(729\\) num_teams\\(4223\\)\[\\n\\r\{ \]+#pragma omp teams thread_limit\\(729\\) num_teams\\(4223\\) message\\(\"16\" \\\[len:2\\\]\\) severity\\(fatal\\)" 1 "original" } }
+
+ !$omp target
+ !$omp teams loop num_teams(444) thread_limit(451) dyn_groupprivate( 842) message("17") ! { dg-message "sorry, unimplemented: 'dyn_groupprivate' clause" }
+ do i = 1, n
+ end do
+ !$omp end target
+! { dg-final { scan-tree-dump-times "#pragma omp target dyn_groupprivate\\(842\\) thread_limit\\(451\\) num_teams\\(444\\)\[\\n\\r\{ \]+#pragma omp teams thread_limit\\(451\\) num_teams\\(444\\) message\\(\"17\" \\\[len:2\\\]\\) severity\\(fatal\\)" 1 "original" } }
+
+
+ ! Same but in two directives, put on target, except for num_teams
+
+ !$omp target thread_limit(strict : 331) dyn_groupprivate( cgroup ,fallback(null) : 243) message("18") ! { dg-message "sorry, unimplemented: 'dyn_groupprivate' clause" }
+ !$omp teams distribute parallel do simd num_teams(11:22) severity(warning)
+ do i = 1, n
+ end do
+ !$omp end target
+! { dg-final { scan-tree-dump-times "#pragma omp target message\\(\"18\" \\\[len:2\\\]\\) severity\\(fatal\\) dyn_groupprivate\\(fallback\\(null\\):243\\) thread_limit\\(strict:331\\) num_teams\\(11:22\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp teams num_teams\\(11:22\\)" 1 "original" } }
+
+ !$omp target thread_limit(relaxed : 1452) dyn_groupprivate( cgroup ,fallback(default_mem) : 1248) message("19") ! { dg-message "sorry, unimplemented: 'dyn_groupprivate' clause" }
+ !$omp teams distribute parallel do num_teams(dims(1):1343) severity(warning)
+ do i = 1, n
+ end do
+ !$omp end target
+! { dg-final { scan-tree-dump-times "#pragma omp target message\\(\"19\" \\\[len:2\\\]\\) severity\\(fatal\\) dyn_groupprivate\\(fallback\\(default_mem\\):1248\\) thread_limit\\(1452\\) num_teams\\(dims\\(\\):1343\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp teams num_teams\\(dims\\(\\):1343\\)" 1 "original" } }
+
+ !$omp target thread_limit(938) dyn_groupprivate( 1458) message("20") ! { dg-message "sorry, unimplemented: 'dyn_groupprivate' clause" }
+ !$omp teams distribute simd num_teams(dims(2):3,3)
+ do i = 1, n
+ end do
+ !$omp end target
+! { dg-final { scan-tree-dump-times "#pragma omp target message\\(\"20\" \\\[len:2\\\]\\) severity\\(fatal\\) dyn_groupprivate\\(1458\\) thread_limit\\(938\\) num_teams\\(dims\\(\\):3, 3\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp teams num_teams\\(dims\\(\\):3, 3\\)" 1 "original" } }
+
+ !$omp target thread_limit(729) dyn_groupprivate( 803) message("21") ! { dg-message "sorry, unimplemented: 'dyn_groupprivate' clause" }
+ !$omp teams distribute num_teams(4223)
+ do i = 1, n
+ end do
+ !$omp end target
+! { dg-final { scan-tree-dump-times "#pragma omp target message\\(\"21\" \\\[len:2\\\]\\) severity\\(fatal\\) dyn_groupprivate\\(803\\) thread_limit\\(729\\) num_teams\\(4223\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp teams num_teams\\(4223\\)" 1 "original" } }
+
+ !$omp target thread_limit(451) dyn_groupprivate( 842) message("22") ! { dg-message "sorry, unimplemented: 'dyn_groupprivate' clause" }
+ !$omp teams loop num_teams(444)
+ do i = 1, n
+ end do
+ !$omp end target
+! { dg-final { scan-tree-dump-times "#pragma omp target message\\(\"22\" \\\[len:2\\\]\\) severity\\(fatal\\) dyn_groupprivate\\(842\\) thread_limit\\(451\\) num_teams\\(444\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp teams num_teams\\(444\\)" 1 "original" } }
+
+
+ ! No 'target', just teams
+
+ !$omp teams distribute parallel do simd num_teams(11:22) thread_limit(strict : 331) dyn_groupprivate( cgroup ,fallback(null) : 243) message("23") severity(warning) ! { dg-message "sorry, unimplemented: 'dyn_groupprivate' clause" }
+ do i = 1, n
+ end do
+! { dg-final { scan-tree-dump-times "#pragma omp teams message\\(\"23\" \\\[len:2\\\]\\) severity\\(fatal\\) dyn_groupprivate\\(fallback\\(null\\):243\\) num_teams\\(11:22\\) thread_limit\\(strict:331\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp parallel message\\(\"23\" \\\[len:2\\\]\\) severity\\(warning\\)" 1 "original" } }
+
+ !$omp teams distribute parallel do num_teams(11:22) thread_limit(strict : 331) dyn_groupprivate( cgroup ,fallback(null) : 243) message("24") severity(warning) ! { dg-message "sorry, unimplemented: 'dyn_groupprivate' clause" }
+ do i = 1, n
+ end do
+! { dg-final { scan-tree-dump-times "#pragma omp teams message\\(\"24\" \\\[len:2\\\]\\) severity\\(fatal\\) dyn_groupprivate\\(fallback\\(null\\):243\\) num_teams\\(11:22\\) thread_limit\\(strict:331\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp parallel message\\(\"24\" \\\[len:2\\\]\\) severity\\(warning\\)" 1 "original" } }
+
+ !$omp teams distribute simd num_teams(11:22) thread_limit(strict : 331) dyn_groupprivate( cgroup ,fallback(null) : 243) message("25") ! { dg-message "sorry, unimplemented: 'dyn_groupprivate' clause" }
+ do i = 1, n
+ end do
+! { dg-final { scan-tree-dump-times "#pragma omp teams message\\(\"25\" \\\[len:2\\\]\\) severity\\(fatal\\) dyn_groupprivate\\(fallback\\(null\\):243\\) num_teams\\(11:22\\) thread_limit\\(strict:331\\)" 1 "original" } }
+
+ !$omp teams distribute num_teams(11:22) thread_limit(strict : 331) dyn_groupprivate( cgroup ,fallback(null) : 243) message("26") ! { dg-message "sorry, unimplemented: 'dyn_groupprivate' clause" }
+ do i = 1, n
+ end do
+! { dg-final { scan-tree-dump-times "#pragma omp teams message\\(\"26\" \\\[len:2\\\]\\) severity\\(fatal\\) dyn_groupprivate\\(fallback\\(null\\):243\\) num_teams\\(11:22\\) thread_limit\\(strict:331\\)" 1 "original" } }
+
+ !$omp teams loop num_teams(11:22) thread_limit(strict : 331) dyn_groupprivate( cgroup ,fallback(null) : 243) message("27") ! { dg-message "sorry, unimplemented: 'dyn_groupprivate' clause" }
+ do i = 1, n
+ end do
+! { dg-final { scan-tree-dump-times "#pragma omp teams message\\(\"27\" \\\[len:2\\\]\\) severity\\(fatal\\) dyn_groupprivate\\(fallback\\(null\\):243\\) num_teams\\(11:22\\) thread_limit\\(strict:331\\)" 1 "original" } }
+end
diff --git a/gcc/testsuite/gfortran.dg/gomp/spatial-dimensions-1.f90 b/gcc/testsuite/gfortran.dg/gomp/spatial-dimensions-1.f90
index 6e000cb49c4..83d3201a215 100644
--- a/gcc/testsuite/gfortran.dg/gomp/spatial-dimensions-1.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/spatial-dimensions-1.f90
@@ -165,7 +165,7 @@ end subroutine
! { dg-final { scan-tree-dump-times "#pragma omp parallel num_threads\\(strict:2\\) shared\\(x\\)" 1 "gimple" } }
! { dg-final { scan-tree-dump-times "#pragma omp teams num_teams\\(1234\\) thread_limit\\(2345\\) shared\\(x\\)" 1 "gimple" } }
! { dg-final { scan-tree-dump-times "#pragma omp parallel num_threads\\(3456\\) shared\\(x\\)" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "#pragma omp teams num_teams\\(123, 456\\) shared\\(x\\)" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "#pragma omp teams num_teams\\(123:456\\) shared\\(x\\)" 1 "gimple" } }
! { dg-final { scan-tree-dump-times "#pragma omp teams num_teams\\(4567\\) thread_limit\\(5678\\) shared\\(x\\)" 1 "gimple" } }
! { dg-final { scan-tree-dump-times "#pragma omp parallel num_threads\\(6789\\) shared\\(x\\)" 1 "gimple" } }
! { dg-final { scan-tree-dump-times "#pragma omp parallel num_threads\\(1, 2, 3\\) shared\\(x\\)" 1 "gimple" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/spatial-dimensions-4.f90 b/gcc/testsuite/gfortran.dg/gomp/spatial-dimensions-4.f90
index 3abca5248f3..059ecb5976d 100644
--- a/gcc/testsuite/gfortran.dg/gomp/spatial-dimensions-4.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/spatial-dimensions-4.f90
@@ -32,17 +32,16 @@ end do
block
x = 1
end block
-! { dg-final { scan-tree-dump-times "#pragma omp target num_teams\\(-1\\) thread_limit\\(31\\) firstprivate\\(x\\) \\\[child fn" 1 "omplower" } }
-! { dg-final { scan-tree-dump-times "#pragma omp teams num_teams\\(dims\\(\\):33\\) thread_limit\\(31\\) shared\\(x\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target thread_limit\\(31\\) num_teams\\(dims\\(\\):33\\) firstprivate\\(x\\) \\\[child fn" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "#pragma omp teams thread_limit\\(31\\) num_teams\\(dims\\(\\):33\\) shared\\(x\\)" 1 "omplower" } }
! { dg-final { scan-tree-dump-times "__builtin_GOMP_teams4 \\(33, 33, 31, D.\[0-9]+\\);" 1 "ompexp" } }
-
!$omp target teams distribute parallel do simd thread_limit(relaxed : 41) num_teams(dims(1) : 43) num_threads( dims ( 1 ) , relaxed : 44)
do i = 1, 1
x = 1
end do
-! { dg-final { scan-tree-dump-times "#pragma omp target num_teams\\(-1\\) thread_limit\\(41\\) map\\(tofrom:i \\\[len: 4\\\] \\\[runtime_implicit\\\]\\) firstprivate\\(x\\) \\\[child fn" 1 "omplower" } }
-! { dg-final { scan-tree-dump-times "#pragma omp teams num_teams\\(dims\\(\\):43\\) thread_limit\\(41\\) shared\\(i\\) shared\\(x\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target thread_limit\\(41\\) num_teams\\(dims\\(\\):43\\) map\\(tofrom:i \\\[len: 4\\\] \\\[runtime_implicit\\\]\\) firstprivate\\(x\\) \\\[child fn" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "#pragma omp teams thread_limit\\(41\\) num_teams\\(dims\\(\\):43\\) shared\\(i\\) shared\\(x\\)" 1 "omplower" } }
! { dg-final { scan-tree-dump-times "#pragma omp parallel _looptemp_\\(D.\[0-9\]+\\) _looptemp_\\(D.\[0-9\]+\\) num_threads\\(dims\\(\\):44\\) lastprivate\\(i\\) shared\\(x\\) \\\[child fn" 1 "omplower" } }
! { dg-final { scan-tree-dump-times "__builtin_GOMP_parallel \\(sub_._omp_fn.\[0-9]+, &.omp_data_o.\[0-9]+, 44, 0\\);" 1 "ompexp" } }
@@ -80,15 +79,15 @@ end block
block
x = 1
end block
-! { dg-final { scan-tree-dump-times "#pragma omp teams num_teams\\(1, 93\\) thread_limit\\(91\\) shared\\(x\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "#pragma omp teams num_teams\\(1:93\\) thread_limit\\(91\\) shared\\(x\\)" 1 "omplower" } }
! { dg-final { scan-tree-dump-times "__builtin_GOMP_teams_reg \\(sub_._omp_fn.\[0-9\]+, &.omp_data_o.\[0-9\]+, 93, 91, 0\\);" 1 "ompexp" } }
!$omp target teams thread_limit(relaxed : 991) num_teams(191 : 993)
block
x = 1
end block
-! { dg-final { scan-tree-dump-times "#pragma omp target num_teams\\(-1\\) thread_limit\\(991\\) firstprivate\\(x\\) \\\[child fn" 1 "omplower" } }
-! { dg-final { scan-tree-dump-times "#pragma omp teams num_teams\\(191, 993\\) thread_limit\\(991\\) shared\\(x\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target thread_limit\\(991\\) num_teams\\(191:993\\) firstprivate\\(x\\) \\\[child fn" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "#pragma omp teams thread_limit\\(991\\) num_teams\\(191:993\\) shared\\(x\\)" 1 "omplower" } }
!$omp teams thread_limit(relaxed : 1231) message("my1234")
block
diff --git a/gcc/testsuite/gfortran.dg/gomp/spatial-dimensions-5.f90 b/gcc/testsuite/gfortran.dg/gomp/spatial-dimensions-5.f90
index 9cde28a76ef..f43678f8ae0 100644
--- a/gcc/testsuite/gfortran.dg/gomp/spatial-dimensions-5.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/spatial-dimensions-5.f90
@@ -9,8 +9,8 @@ integer :: x, i
do i = 1,1
x = 0
end do
-! { dg-final { scan-tree-dump-times "#pragma omp target num_teams\\(-1\\) message\\(.first. \\\[len:5\\\]\\) severity\\(fatal\\) thread_limit\\(strict,dims\\(\\):11, 12, 13\\) firstprivate\\(x\\)" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "#pragma omp teams num_teams\\(11, 22\\) thread_limit\\(strict,dims\\(\\):11, 12, 13\\) message\\(.first. \\\[len:5\\\]\\) severity\\(fatal\\) shared\\(x\\)" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target message\\(.first. \\\[len:5\\\]\\) severity\\(fatal\\) thread_limit\\(strict,dims\\(\\):11, 12, 13\\) num_teams\\(11:22\\) firstprivate\\(x\\)" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "#pragma omp teams thread_limit\\(strict,dims\\(\\):11, 12, 13\\) num_teams\\(11:22\\) message\\(.first. \\\[len:5\\\]\\) severity\\(fatal\\) shared\\(x\\)" 1 "gimple" } }
! { dg-final { scan-tree-dump-times "#pragma omp parallel message\\(.first. \\\[len:5\\\]\\) severity\\(warning\\) num_threads\\(23, 45, 53\\) shared\\(x\\)" 1 "gimple" } }
end
@@ -27,15 +27,15 @@ integer :: x, i
do i = 1,1
x = 0
end do
-! { dg-final { scan-tree-dump-times "#pragma omp target num_teams\\(-1\\) message\\(.second. \\\[len:6\\\]\\) severity\\(fatal\\) thread_limit\\(11\\) firstprivate\\(x\\)" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "#pragma omp teams num_teams\\(dims\\(\\):111, 222, 333, 324\\) thread_limit\\(11\\) message\\(.second. \\\[len:6\\\]\\) severity\\(fatal\\) shared\\(x\\)" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target message\\(.second. \\\[len:6\\\]\\) severity\\(fatal\\) thread_limit\\(11\\) num_teams\\(dims\\(\\):111, 222, 333, 324\\) firstprivate\\(x\\)" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "#pragma omp teams thread_limit\\(11\\) num_teams\\(dims\\(\\):111, 222, 333, 324\\) message\\(.second. \\\[len:6\\\]\\) severity\\(fatal\\) shared\\(x\\)" 1 "gimple" } }
! { dg-final { scan-tree-dump-times "#pragma omp parallel message\\(.second. \\\[len:6\\\]\\) severity\\(fatal\\) num_threads\\(strict:445, 553\\) shared\\(x\\)" 1 "gimple" } }
!$omp teams distribute parallel do num_teams(11 : 55) thread_limit(strict : 11) num_threads(relaxed: 23,45,53) message("third") severity(warning)
do i = 1,1
x = 0
end do
-! { dg-final { scan-tree-dump-times "#pragma omp teams message\\(.third. \\\[len:5\\\]\\) severity\\(fatal\\) num_teams\\(11, 55\\) thread_limit\\(strict:11\\) shared\\(x\\)" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "#pragma omp teams message\\(.third. \\\[len:5\\\]\\) severity\\(fatal\\) num_teams\\(11:55\\) thread_limit\\(strict:11\\) shared\\(x\\)" 1 "gimple" } }
! { dg-final { scan-tree-dump-times "#pragma omp parallel message\\(.third. \\\[len:5\\\]\\) severity\\(warning\\) num_threads\\(23, 45, 53\\) shared\\(x\\)" 1 "gimple" } }
!$omp teams distribute parallel do num_teams(dims(2): 1111,33) thread_limit(dims(3) : 11,3,480) num_threads(strict, dims(2): 445,553) message("fourth") severity(fatal)
@@ -61,7 +61,7 @@ end block
block
x = 0
end block
-! { dg-final { scan-tree-dump-times "#pragma omp teams message\\(.teams. \\\[len:5\\\]\\) severity\\(fatal\\) num_teams\\(14, 34\\) shared\\(x\\)" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "#pragma omp teams message\\(.teams. \\\[len:5\\\]\\) severity\\(fatal\\) num_teams\\(14:34\\) shared\\(x\\)" 1 "gimple" } }
!$omp teams thread_limit(relaxed, dims(2): 514 , 384)
block
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-teams-1.f90 b/gcc/testsuite/gfortran.dg/gomp/target-teams-1.f90
new file mode 100644
index 00000000000..39387a9878a
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/target-teams-1.f90
@@ -0,0 +1,65 @@
+! { dg-additional-options "-fdump-tree-omplower" }
+
+subroutine foo
+ implicit none
+
+ !$omp target teams num_teams(dims(2) : 1,2) ! { dg-message "sorry, unimplemented: 'num_teams' clause with 'dims' modifier" }
+ !$omp end target teams
+
+ !$omp target
+ !$omp teams num_teams(dims(2) : 2,3) ! { dg-message "sorry, unimplemented: 'num_teams' clause with 'dims' modifier" }
+ !$omp end teams
+ !$omp end target
+
+ !$omp teams num_teams(dims(2) : 4,5) ! { dg-message "sorry, unimplemented: 'num_teams' clause with 'dims' modifier" }
+ !$omp end teams
+
+ !$omp target teams num_teams(dims(2) : 5,6) ! { dg-message "sorry, unimplemented: 'num_teams' clause with 'dims' modifier" }
+ !$omp end target teams
+
+ ! -----
+
+ !$omp target thread_limit(dims(3), relaxed : 1,2,3) ! { dg-message "sorry, unimplemented: 'thread_limit' clause with 'dims' modifier" }
+ !$omp end target
+
+ !$omp target thread_limit(dims(3), strict : 2,3,4) ! { dg-message "sorry, unimplemented: 'thread_limit' clause with 'dims' modifier" }
+ !$omp teams
+ !$omp end teams
+ !$omp end target
+
+ !$omp target teams thread_limit(dims(3), strict : 3,4,5) ! { dg-message "sorry, unimplemented: 'thread_limit' clause with 'dims' modifier" }
+ !$omp end target teams
+
+ !$omp target
+ !$omp teams thread_limit(strict : 456) ! { dg-message "sorry, unimplemented: 'thread_limit' clause with 'strict' modifier" }
+ !$omp end teams
+ !$omp end target
+
+ !$omp teams thread_limit(strict : 567) ! { dg-message "sorry, unimplemented: 'thread_limit' clause with 'strict' modifier" }
+ !$omp end teams
+end
+
+! { dg-final { scan-tree-dump-times "#pragma omp target num_teams\\(dims\\(\\):1, 2\\) \\\[child fn:" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "#pragma omp teams num_teams\\(dims\\(\\):1, 2\\)" 1 "omplower" } }
+
+! { dg-final { scan-tree-dump-times "#pragma omp target num_teams\\(dims\\(\\):2, 3\\) \\\[child fn:" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "#pragma omp teams num_teams\\(dims\\(\\):2, 3\\)" 1 "omplower" } }
+
+! { dg-final { scan-tree-dump-times "#pragma omp teams num_teams\\(dims\\(\\):4, 5\\)" 1 "omplower" } }
+
+! { dg-final { scan-tree-dump-times "#pragma omp target num_teams\\(dims\\(\\):5, 6\\) \\\[child fn:" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "#pragma omp teams num_teams\\(dims\\(\\):5, 6\\)" 1 "omplower" } }
+
+! { dg-final { scan-tree-dump-times "#pragma omp target num_teams\\(-2\\) thread_limit\\(dims\\(\\):1, 2, 3\\) \\\[child fn:" 1 "omplower" } }
+
+! { dg-final { scan-tree-dump-times "#pragma omp target num_teams\\(0\\) thread_limit\\(strict,dims\\(\\):2, 3, 4\\) \\\[child fn:" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "#pragma omp teams\[\\n\\r\]" 1 "omplower" } }
+
+! { dg-final { scan-tree-dump-times "#pragma omp target num_teams\\(0\\) thread_limit\\(strict,dims\\(\\):3, 4, 5\\) \\\[child fn:" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "#pragma omp teams thread_limit\\(strict,dims\\(\\):3, 4, 5\\)" 1 "omplower" } }
+
+! { dg-final { scan-tree-dump-times "#pragma omp target num_teams\\(0\\) thread_limit\\(strict:456\\) \\\[child fn:" 1 "omplower" } }
+
+! { dg-final { scan-tree-dump-times "#pragma omp teams thread_limit\\(strict:456\\)" 1 "omplower" } }
+
+! { dg-final { scan-tree-dump-times "#pragma omp teams thread_limit\\(strict:567\\)" 1 "omplower" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-teams-2.f90 b/gcc/testsuite/gfortran.dg/gomp/target-teams-2.f90
new file mode 100644
index 00000000000..8098f5c82df
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/target-teams-2.f90
@@ -0,0 +1,17 @@
+! { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" }
+
+! Check that the expression is evaluated before the target region
+
+subroutine sub(n)
+implicit none
+integer, value :: n
+
+!$omp target
+!$omp teams thread_limit(n+1)
+!$omp end teams
+!$omp end target
+end
+! { dg-final { scan-tree-dump "D\.\[0-9\]+ = n \\+ 1;\[\\n\\r\]+ *#pragma omp target thread_limit\\(D\.\[0-9\]+\\) firstprivate\\(D\.\[0-9\]+\\)\[\\n\\r\{ \]+#pragma omp teams thread_limit\\(D\.\[0-9\]+\\)" "original" } }
+
+! { dg-final { scan-tree-dump "D\.\[0-9\]+ = n \\+ 1;\[\\n\\r\]+ *_\[0-9\]+ = D\.\[0-9\]+;\[\\n\\r\]+ *#pragma omp target num_teams\\(0\\) thread_limit\\(_\[0-9\]+\\) firstprivate\\(D\.\[0-9\]+\\)\[\\n\\r\{ \]+integer\\(kind=4\\) D\.\[0-9\]+;\[\\n\\r \]+D\.\[0-9\]+ = D\.\[0-9\]+;\[\\n\\r \]+#pragma omp teams thread_limit\\(D\.\[0-9\]+\\)" "gimple" } }
+
