This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
[gomp4.5] target teams expression evaluation
- From: Jakub Jelinek <jakub at redhat dot com>
- To: gcc-patches at gcc dot gnu dot org
- Cc: Alexander Monakov <amonakov at ispras dot ru>, Ilya Verbin <iverbin at gmail dot com>
- Date: Mon, 26 Oct 2015 15:51:17 +0100
- Subject: [gomp4.5] target teams expression evaluation
- Authentication-results: sourceware.org; auth=none
- Reply-to: Jakub Jelinek <jakub at redhat dot com>
Hi!
The OpenMP 4.5 spec says that for combined target teams the num_teams
and thread_limit expressions are evaluated on the host before the
target construct.
Additionally, this patch tries to determine during gimplification if
it is easily possible to compute those expressions on the host even
if the construct is not combined (basically if the expressions are simple
integral arithmetics using constants and variables that are either
explicitly firstprivate, or explicitly mapped always {to,tofrom}, or
implicitly firstprivatized on the target construct.
The runtime library function is then told these values (> 0 stands
for those values, 0 stands for clauses not specified, meaning implementation
defined values can be used, -1 stands for could not determine what will be
passed; for missing teams construct which is essentially a request to have
exactly 1 team and implementation defined number of maximum threads we pass
1, 0).
Regtested on x86_64-linux, both intelmicemul and no offloading.
As incremental change, probably the GOMP_target_ext construct could serve
as GOMP_teams call too if both of those values are >= 0 and then GOMP_teams
call should be removed.
2015-10-26 Jakub Jelinek <jakub@redhat.com>
gcc/
* gimplify.c (enum gimplify_omp_var_data): Add GOVD_MAP_ALWAYS_TO.
(gimplify_scan_omp_clauses): Or in GOVD_MAP_ALWAYS_TO for
GOMP_MAP_ALWAYS_TO or GOMP_MAP_ALWAYS_TOFROM kinds.
(find_omp_teams, computable_teams_clause, optimize_target_teams): New
functions.
(gimplify_omp_workshare): Call optimize_target_teams.
* omp-low.c (expand_omp_target): Pass num_teams and thread_limit
arguments to BUILT_IN_GOMP_TARGET.
* omp-builtins.def (BUILT_IN_GOMP_TARGET): Rename GOMP_target_41
to GOMP_target_ext. Add num_teams and thread_limit arguments.
(BUILT_IN_GOMP_TARGET_DATA): Rename GOMP_target_data_41
to GOMP_target_data_ext.
(BUILT_IN_GOMP_TARGET_UPDATE): Rename GOMP_target_update_41
to GOMP_target_update_ext.
* builtin-types.def
(BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR): Remove.
(BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR_INT_INT): New.
gcc/c/
* c-parser.c: Include gimple-expr.h.
(c_parser_omp_target): Evaluate num_teams and thread_limit
expressions on combined target teams before the target.
gcc/cp/
* parser.c (cp_parser_omp_target): Evaluate num_teams and
thread_limit expressions on combined target teams before the
target.
* pt.c (tsubst_find_omp_teams): New function.
(tsubst_expr): Evaluate num_teams and thread_limit expressions on
combined target teams before the target.
gcc/fortran/
* types.def (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR): Remove.
(BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR_INT_INT): New.
gcc/testsuite/
* c-c++-common/gomp/target-teams-1.c: New test.
* g++.dg/gomp/target-teams-1.C: New test.
include/
* gomp-constants.h (GOMP_TARGET_FLAG_NOWAIT): Adjust comment.
libgomp/
* target.c (GOMP_target_41): Renamed to ...
(GOMP_target_ext): ... this. Add num_teams and thread_limit
arguments.
(GOMP_target_data_41): Renamed to ...
(GOMP_target_data_ext): ... this.
(GOMP_target_update_41): Renamed to ...
(GOMP_target_update_ext): ... this.
* libgomp.map (GOMP_4.5): Export GOMP_target_ext,
GOMP_target_data_ext and GOMP_target_update_ext instead of
GOMP_target_41, GOMP_target_data_41 and GOMP_target_update_41.
* libgomp_g.h (GOMP_target_41): Renamed to ...
(GOMP_target_ext): ... this. Add num_teams and thread_limit
arguments.
(GOMP_target_data_41): Renamed to ...
(GOMP_target_data_ext): ... this.
(GOMP_target_update_41): Renamed to ...
(GOMP_target_update_ext): ... this.
* testsuite/libgomp.c/target-teams-1.c: New test.
--- gcc/gimplify.c.jj 2015-10-22 18:01:45.000000000 +0200
+++ gcc/gimplify.c 2015-10-23 17:44:44.891528688 +0200
@@ -93,6 +93,9 @@ enum gimplify_omp_var_data
GOVD_MAP_0LEN_ARRAY = 32768,
+ /* Flag for GOVD_MAP, if it is always, to or always, tofrom mapping. */
+ GOVD_MAP_ALWAYS_TO = 65536,
+
GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
| GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
| GOVD_LOCAL)
@@ -6757,6 +6760,9 @@ gimplify_scan_omp_clauses (tree *list_p,
break;
}
flags = GOVD_MAP | GOVD_EXPLICIT;
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TO
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TOFROM)
+ flags |= GOVD_MAP_ALWAYS_TO;
goto do_add;
case OMP_CLAUSE_DEPEND:
@@ -8543,6 +8549,201 @@ gimplify_omp_for (tree *expr_p, gimple_s
return GS_ALL_DONE;
}
+/* Helper function of optimize_target_teams, find OMP_TEAMS inside
+ of OMP_TARGET's body. */
+
+static tree
+find_omp_teams (tree *tp, int *walk_subtrees, void *)
+{
+ *walk_subtrees = 0;
+ switch (TREE_CODE (*tp))
+ {
+ case OMP_TEAMS:
+ return *tp;
+ case BIND_EXPR:
+ case STATEMENT_LIST:
+ *walk_subtrees = 1;
+ break;
+ default:
+ break;
+ }
+ return NULL_TREE;
+}
+
+/* Helper function of optimize_target_teams, determine if the expression
+ can be computed safely before the target construct on the host. */
+
+static tree
+computable_teams_clause (tree *tp, int *walk_subtrees, void *)
+{
+ splay_tree_node n;
+
+ if (TYPE_P (*tp))
+ {
+ *walk_subtrees = 0;
+ return NULL_TREE;
+ }
+ switch (TREE_CODE (*tp))
+ {
+ case VAR_DECL:
+ case PARM_DECL:
+ case RESULT_DECL:
+ *walk_subtrees = 0;
+ if (error_operand_p (*tp)
+ || !INTEGRAL_TYPE_P (TREE_TYPE (*tp))
+ || DECL_HAS_VALUE_EXPR_P (*tp)
+ || DECL_THREAD_LOCAL_P (*tp)
+ || TREE_SIDE_EFFECTS (*tp)
+ || TREE_THIS_VOLATILE (*tp))
+ return *tp;
+ if (is_global_var (*tp)
+ && (lookup_attribute ("omp declare target", DECL_ATTRIBUTES (*tp))
+ || lookup_attribute ("omp declare target link",
+ DECL_ATTRIBUTES (*tp))))
+ return *tp;
+ n = splay_tree_lookup (gimplify_omp_ctxp->variables,
+ (splay_tree_key) *tp);
+ if (n == NULL)
+ {
+ if (gimplify_omp_ctxp->target_map_scalars_firstprivate)
+ return NULL_TREE;
+ return *tp;
+ }
+ else if (n->value & GOVD_LOCAL)
+ return *tp;
+ else if (n->value & GOVD_FIRSTPRIVATE)
+ return NULL_TREE;
+ else if ((n->value & (GOVD_MAP | GOVD_MAP_ALWAYS_TO))
+ == (GOVD_MAP | GOVD_MAP_ALWAYS_TO))
+ return NULL_TREE;
+ return *tp;
+ case INTEGER_CST:
+ if (!INTEGRAL_TYPE_P (TREE_TYPE (*tp)))
+ return *tp;
+ return NULL_TREE;
+ case TARGET_EXPR:
+ if (TARGET_EXPR_INITIAL (*tp)
+ || TREE_CODE (TARGET_EXPR_SLOT (*tp)) != VAR_DECL)
+ return *tp;
+ return computable_teams_clause (&TARGET_EXPR_SLOT (*tp),
+ walk_subtrees, NULL);
+ /* Allow some reasonable subset of integral arithmetics. */
+ case PLUS_EXPR:
+ case MINUS_EXPR:
+ case MULT_EXPR:
+ case TRUNC_DIV_EXPR:
+ case CEIL_DIV_EXPR:
+ case FLOOR_DIV_EXPR:
+ case ROUND_DIV_EXPR:
+ case TRUNC_MOD_EXPR:
+ case CEIL_MOD_EXPR:
+ case FLOOR_MOD_EXPR:
+ case ROUND_MOD_EXPR:
+ case RDIV_EXPR:
+ case EXACT_DIV_EXPR:
+ case MIN_EXPR:
+ case MAX_EXPR:
+ case LSHIFT_EXPR:
+ case RSHIFT_EXPR:
+ case BIT_IOR_EXPR:
+ case BIT_XOR_EXPR:
+ case BIT_AND_EXPR:
+ case NEGATE_EXPR:
+ case ABS_EXPR:
+ case BIT_NOT_EXPR:
+ case NON_LVALUE_EXPR:
+ CASE_CONVERT:
+ if (!INTEGRAL_TYPE_P (TREE_TYPE (*tp)))
+ return *tp;
+ return NULL_TREE;
+ /* And disallow anything else, except for comparisons. */
+ default:
+ if (COMPARISON_CLASS_P (*tp))
+ return NULL_TREE;
+ return *tp;
+ }
+}
+
+/* Try to determine if the num_teams and/or thread_limit expressions
+ can have their values determined already before entering the
+ target construct.
+ INTEGER_CSTs trivially are,
+ integral decls that are firstprivate (explicitly or implicitly)
+ or explicitly map(always, to:) or map(always, tofrom:) on the target
+ region too, and expressions involving simple arithmetics on those
+ too, function calls are not ok, dereferencing something neither etc.
+ Add NUM_TEAMS and THREAD_LIMIT clauses to the OMP_CLAUSES of
+ EXPR based on what we find:
+ 0 stands for clause not specified at all, use implementation default
+ -1 stands for value that can't be determined easily before entering
+ the target construct.
+ If teams construct is not present at all, use 1 for num_teams
+ and 0 for thread_limit (only one team is involved, and the thread
+ limit is implementation defined. */
+
+static void
+optimize_target_teams (tree target, gimple_seq *pre_p)
+{
+ tree body = OMP_BODY (target);
+ tree teams = walk_tree (&body, find_omp_teams, NULL, NULL);
+ tree num_teams = integer_zero_node;
+ tree thread_limit = integer_zero_node;
+ location_t num_teams_loc = EXPR_LOCATION (target);
+ location_t thread_limit_loc = EXPR_LOCATION (target);
+ tree c, *p, expr;
+ struct gimplify_omp_ctx *target_ctx = gimplify_omp_ctxp;
+
+ if (teams == NULL_TREE)
+ num_teams = integer_one_node;
+ else
+ for (c = OMP_TEAMS_CLAUSES (teams); c; c = OMP_CLAUSE_CHAIN (c))
+ {
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_NUM_TEAMS)
+ {
+ p = &num_teams;
+ num_teams_loc = OMP_CLAUSE_LOCATION (c);
+ }
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_THREAD_LIMIT)
+ {
+ p = &thread_limit;
+ thread_limit_loc = OMP_CLAUSE_LOCATION (c);
+ }
+ else
+ continue;
+ expr = OMP_CLAUSE_OPERAND (c, 0);
+ if (TREE_CODE (expr) == INTEGER_CST)
+ {
+ *p = expr;
+ continue;
+ }
+ if (walk_tree (&expr, computable_teams_clause, NULL, NULL))
+ {
+ *p = integer_minus_one_node;
+ continue;
+ }
+ *p = expr;
+ gimplify_omp_ctxp = gimplify_omp_ctxp->outer_context;
+ if (gimplify_expr (p, pre_p, NULL, is_gimple_val, fb_rvalue)
+ == GS_ERROR)
+ {
+ gimplify_omp_ctxp = target_ctx;
+ *p = integer_minus_one_node;
+ continue;
+ }
+ gimplify_omp_ctxp = target_ctx;
+ if (!DECL_P (expr) && TREE_CODE (expr) != TARGET_EXPR)
+ OMP_CLAUSE_OPERAND (c, 0) = *p;
+ }
+ c = build_omp_clause (thread_limit_loc, OMP_CLAUSE_THREAD_LIMIT);
+ OMP_CLAUSE_THREAD_LIMIT_EXPR (c) = thread_limit;
+ OMP_CLAUSE_CHAIN (c) = OMP_TARGET_CLAUSES (target);
+ OMP_TARGET_CLAUSES (target) = c;
+ c = build_omp_clause (num_teams_loc, OMP_CLAUSE_NUM_TEAMS);
+ OMP_CLAUSE_NUM_TEAMS_EXPR (c) = num_teams;
+ OMP_CLAUSE_CHAIN (c) = OMP_TARGET_CLAUSES (target);
+ OMP_TARGET_CLAUSES (target) = c;
+}
+
/* Gimplify the gross structure of several OMP constructs. */
static void
@@ -8578,6 +8779,8 @@ gimplify_omp_workshare (tree *expr_p, gi
}
gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ort,
TREE_CODE (expr));
+ if (TREE_CODE (expr) == OMP_TARGET)
+ optimize_target_teams (expr, pre_p);
if ((ort & (ORT_TARGET | ORT_TARGET_DATA)) != 0)
{
push_gimplify_context ();
--- gcc/omp-low.c.jj 2015-10-23 13:10:16.000000000 +0200
+++ gcc/omp-low.c 2015-10-26 13:14:33.760665067 +0100
@@ -11857,6 +11857,31 @@ expand_omp_target (struct omp_region *re
else
depend = build_int_cst (ptr_type_node, 0);
args.quick_push (depend);
+ if (start_ix == BUILT_IN_GOMP_TARGET)
+ {
+ c = find_omp_clause (clauses, OMP_CLAUSE_NUM_TEAMS);
+ if (c)
+ {
+ t = fold_convert (integer_type_node,
+ OMP_CLAUSE_NUM_TEAMS_EXPR (c));
+ t = force_gimple_operand_gsi (&gsi, t, true, NULL,
+ true, GSI_SAME_STMT);
+ }
+ else
+ t = integer_minus_one_node;
+ args.quick_push (t);
+ c = find_omp_clause (clauses, OMP_CLAUSE_THREAD_LIMIT);
+ if (c)
+ {
+ t = fold_convert (integer_type_node,
+ OMP_CLAUSE_THREAD_LIMIT_EXPR (c));
+ t = force_gimple_operand_gsi (&gsi, t, true, NULL,
+ true, GSI_SAME_STMT);
+ }
+ else
+ t = integer_minus_one_node;
+ args.quick_push (t);
+ }
break;
case BUILT_IN_GOACC_PARALLEL:
{
--- gcc/omp-builtins.def.jj 2015-10-14 18:04:13.000000000 +0200
+++ gcc/omp-builtins.def 2015-10-26 13:03:14.345593535 +0100
@@ -342,14 +342,14 @@ DEF_GOMP_BUILTIN (BUILT_IN_GOMP_SINGLE_C
BT_FN_PTR, ATTR_NOTHROW_LEAF_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_GOMP_SINGLE_COPY_END, "GOMP_single_copy_end",
BT_FN_VOID_PTR, ATTR_NOTHROW_LEAF_LIST)
-DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET, "GOMP_target_41",
- BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR,
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET, "GOMP_target_ext",
+ BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR_INT_INT,
ATTR_NOTHROW_LIST)
-DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_DATA, "GOMP_target_data_41",
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_DATA, "GOMP_target_data_ext",
BT_FN_VOID_INT_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_END_DATA, "GOMP_target_end_data",
BT_FN_VOID, ATTR_NOTHROW_LIST)
-DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_UPDATE, "GOMP_target_update_41",
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_UPDATE, "GOMP_target_update_ext",
BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_UINT_PTR,
ATTR_NOTHROW_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_ENTER_EXIT_DATA,
--- gcc/builtin-types.def.jj 2015-10-14 10:25:32.000000000 +0200
+++ gcc/builtin-types.def 2015-10-26 13:04:01.178908933 +0100
@@ -547,9 +547,6 @@ DEF_FUNCTION_TYPE_7 (BT_FN_VOID_INT_SIZE
BT_VOID, BT_INT, BT_SIZE, BT_PTR, BT_PTR, BT_PTR, BT_UINT,
BT_PTR)
-DEF_FUNCTION_TYPE_8 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR,
- BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE, BT_PTR,
- BT_PTR, BT_PTR, BT_UINT, BT_PTR)
DEF_FUNCTION_TYPE_8 (BT_FN_VOID_OMPFN_PTR_UINT_LONG_LONG_LONG_LONG_UINT,
BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR, BT_UINT,
BT_LONG, BT_LONG, BT_LONG, BT_LONG, BT_UINT)
@@ -559,6 +556,10 @@ DEF_FUNCTION_TYPE_9 (BT_FN_VOID_OMPFN_PT
BT_PTR_FN_VOID_PTR_PTR, BT_LONG, BT_LONG,
BT_BOOL, BT_UINT, BT_PTR, BT_INT)
+DEF_FUNCTION_TYPE_10 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR_INT_INT,
+ BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE, BT_PTR,
+ BT_PTR, BT_PTR, BT_UINT, BT_PTR, BT_INT, BT_INT)
+
DEF_FUNCTION_TYPE_11 (BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_UINT_LONG_INT_LONG_LONG_LONG,
BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR,
BT_PTR_FN_VOID_PTR_PTR, BT_LONG, BT_LONG,
--- gcc/c/c-parser.c.jj 2015-10-22 14:35:41.000000000 +0200
+++ gcc/c/c-parser.c 2015-10-23 14:48:25.849665994 +0200
@@ -66,6 +66,7 @@ along with GCC; see the file COPYING3.
#include "builtins.h"
#include "gomp-constants.h"
#include "c-family/c-indentation.h"
+#include "gimple-expr.h"
/* Initialization routine for this file. */
@@ -15202,6 +15203,32 @@ c_parser_omp_target (c_parser *parser, e
block = c_end_compound_stmt (loc, block, true);
if (ret == NULL_TREE)
return false;
+ if (ccode == OMP_TEAMS)
+ {
+ /* For combined target teams, ensure the num_teams and
+ thread_limit clause expressions are evaluated on the host,
+ before entering the target construct. */
+ tree c;
+ for (c = cclauses[C_OMP_CLAUSE_SPLIT_TEAMS];
+ c; c = OMP_CLAUSE_CHAIN (c))
+ if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_NUM_TEAMS
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_THREAD_LIMIT)
+ && TREE_CODE (OMP_CLAUSE_OPERAND (c, 0)) != INTEGER_CST)
+ {
+ tree expr = OMP_CLAUSE_OPERAND (c, 0);
+ tree tmp = create_tmp_var_raw (TREE_TYPE (expr));
+ expr = build4 (TARGET_EXPR, TREE_TYPE (expr), tmp,
+ expr, NULL_TREE, NULL_TREE);
+ add_stmt (expr);
+ OMP_CLAUSE_OPERAND (c, 0) = expr;
+ tree tc = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+ OMP_CLAUSE_FIRSTPRIVATE);
+ OMP_CLAUSE_DECL (tc) = tmp;
+ OMP_CLAUSE_CHAIN (tc)
+ = cclauses[C_OMP_CLAUSE_SPLIT_TARGET];
+ cclauses[C_OMP_CLAUSE_SPLIT_TARGET] = tc;
+ }
+ }
tree stmt = make_node (OMP_TARGET);
TREE_TYPE (stmt) = void_type_node;
OMP_TARGET_CLAUSES (stmt) = cclauses[C_OMP_CLAUSE_SPLIT_TARGET];
--- gcc/cp/parser.c.jj 2015-10-22 14:35:41.000000000 +0200
+++ gcc/cp/parser.c 2015-10-26 11:53:55.469514327 +0100
@@ -34116,6 +34116,33 @@ cp_parser_omp_target (cp_parser *parser,
tree body = finish_omp_structured_block (sb);
if (ret == NULL_TREE)
return false;
+ if (ccode == OMP_TEAMS && !processing_template_decl)
+ {
+ /* For combined target teams, ensure the num_teams and
+ thread_limit clause expressions are evaluated on the host,
+ before entering the target construct. */
+ tree c;
+ for (c = cclauses[C_OMP_CLAUSE_SPLIT_TEAMS];
+ c; c = OMP_CLAUSE_CHAIN (c))
+ if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_NUM_TEAMS
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_THREAD_LIMIT)
+ && TREE_CODE (OMP_CLAUSE_OPERAND (c, 0)) != INTEGER_CST)
+ {
+ tree expr = OMP_CLAUSE_OPERAND (c, 0);
+ expr = force_target_expr (TREE_TYPE (expr), expr, tf_none);
+ if (expr == error_mark_node)
+ continue;
+ tree tmp = TARGET_EXPR_SLOT (expr);
+ add_stmt (expr);
+ OMP_CLAUSE_OPERAND (c, 0) = expr;
+ tree tc = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+ OMP_CLAUSE_FIRSTPRIVATE);
+ OMP_CLAUSE_DECL (tc) = tmp;
+ OMP_CLAUSE_CHAIN (tc)
+ = cclauses[C_OMP_CLAUSE_SPLIT_TARGET];
+ cclauses[C_OMP_CLAUSE_SPLIT_TARGET] = tc;
+ }
+ }
tree stmt = make_node (OMP_TARGET);
TREE_TYPE (stmt) = void_type_node;
OMP_TARGET_CLAUSES (stmt) = cclauses[C_OMP_CLAUSE_SPLIT_TARGET];
--- gcc/cp/pt.c.jj 2015-10-19 14:44:36.000000000 +0200
+++ gcc/cp/pt.c 2015-10-26 12:12:08.171514471 +0100
@@ -14813,6 +14813,27 @@ tsubst_omp_for_iterator (tree t, int i,
#undef RECUR
}
+/* Helper function of tsubst_expr, find OMP_TEAMS inside
+ of OMP_TARGET's body. */
+
+static tree
+tsubst_find_omp_teams (tree *tp, int *walk_subtrees, void *)
+{
+ *walk_subtrees = 0;
+ switch (TREE_CODE (*tp))
+ {
+ case OMP_TEAMS:
+ return *tp;
+ case BIND_EXPR:
+ case STATEMENT_LIST:
+ *walk_subtrees = 1;
+ break;
+ default:
+ break;
+ }
+ return NULL_TREE;
+}
+
/* Like tsubst_copy for expressions, etc. but also does semantic
processing. */
@@ -15315,6 +15336,36 @@ tsubst_expr (tree t, tree args, tsubst_f
t = copy_node (t);
OMP_BODY (t) = stmt;
OMP_CLAUSES (t) = tmp;
+ if (TREE_CODE (t) == OMP_TARGET && OMP_TARGET_COMBINED (t))
+ {
+ tree teams = cp_walk_tree (&stmt, tsubst_find_omp_teams, NULL, NULL);
+ if (teams)
+ {
+ /* For combined target teams, ensure the num_teams and
+ thread_limit clause expressions are evaluated on the host,
+ before entering the target construct. */
+ tree c;
+ for (c = OMP_TEAMS_CLAUSES (teams);
+ c; c = OMP_CLAUSE_CHAIN (c))
+ if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_NUM_TEAMS
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_THREAD_LIMIT)
+ && TREE_CODE (OMP_CLAUSE_OPERAND (c, 0)) != INTEGER_CST)
+ {
+ tree expr = OMP_CLAUSE_OPERAND (c, 0);
+ expr = force_target_expr (TREE_TYPE (expr), expr, tf_none);
+ if (expr == error_mark_node)
+ continue;
+ tmp = TARGET_EXPR_SLOT (expr);
+ add_stmt (expr);
+ OMP_CLAUSE_OPERAND (c, 0) = expr;
+ tree tc = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+ OMP_CLAUSE_FIRSTPRIVATE);
+ OMP_CLAUSE_DECL (tc) = tmp;
+ OMP_CLAUSE_CHAIN (tc) = OMP_TARGET_CLAUSES (t);
+ OMP_TARGET_CLAUSES (t) = tc;
+ }
+ }
+ }
add_stmt (t);
break;
--- gcc/fortran/types.def.jj 2015-10-14 10:25:39.000000000 +0200
+++ gcc/fortran/types.def 2015-10-26 13:04:41.284322679 +0100
@@ -215,15 +215,16 @@ DEF_FUNCTION_TYPE_7 (BT_FN_VOID_INT_SIZE
DEF_FUNCTION_TYPE_8 (BT_FN_VOID_OMPFN_PTR_UINT_LONG_LONG_LONG_LONG_UINT,
BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR, BT_UINT,
BT_LONG, BT_LONG, BT_LONG, BT_LONG, BT_UINT)
-DEF_FUNCTION_TYPE_8 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR,
- BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE, BT_PTR,
- BT_PTR, BT_PTR, BT_UINT, BT_PTR)
DEF_FUNCTION_TYPE_9 (BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR_INT,
BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR,
BT_PTR_FN_VOID_PTR_PTR, BT_LONG, BT_LONG,
BT_BOOL, BT_UINT, BT_PTR, BT_INT)
+DEF_FUNCTION_TYPE_10 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR_INT_INT,
+ BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE, BT_PTR,
+ BT_PTR, BT_PTR, BT_UINT, BT_PTR, BT_INT, BT_INT)
+
DEF_FUNCTION_TYPE_11 (BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_UINT_LONG_INT_LONG_LONG_LONG,
BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR,
BT_PTR_FN_VOID_PTR_PTR, BT_LONG, BT_LONG,
--- gcc/testsuite/c-c++-common/gomp/target-teams-1.c.jj 2015-10-23 18:12:39.466590989 +0200
+++ gcc/testsuite/c-c++-common/gomp/target-teams-1.c 2015-10-23 18:37:02.487625789 +0200
@@ -0,0 +1,85 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp -fdump-tree-gimple" } */
+
+int v = 6;
+void bar (int);
+void bar2 (int, long *, long *);
+int baz (void);
+#pragma omp declare target to (bar, baz, v)
+
+void
+foo (int a, int b, long c, long d)
+{
+ /* The OpenMP 4.5 spec says that these expressions are evaluated before
+ target region on combined target teams, so those cases are always
+ fine. */
+ #pragma omp target
+ bar (0);
+ #pragma omp target
+ #pragma omp teams
+ bar (1);
+ #pragma omp target teams
+ bar (2);
+ #pragma omp target
+ #pragma omp teams num_teams (4)
+ bar (3);
+ #pragma omp target teams num_teams (4)
+ bar (4);
+ #pragma omp target
+ #pragma omp teams thread_limit (7)
+ bar (5);
+ #pragma omp target teams thread_limit (7)
+ bar (6);
+ #pragma omp target
+ #pragma omp teams num_teams (4) thread_limit (8)
+ {
+ {
+ bar (7);
+ }
+ }
+ #pragma omp target teams num_teams (4) thread_limit (8)
+ bar (8);
+ #pragma omp target
+ #pragma omp teams num_teams (a) thread_limit (b)
+ bar (9);
+ #pragma omp target teams num_teams (a) thread_limit (b)
+ bar (10);
+ #pragma omp target
+ #pragma omp teams num_teams (c + 1) thread_limit (d - 1)
+ bar (11);
+ #pragma omp target teams num_teams (c + 1) thread_limit (d - 1)
+ bar (12);
+ #pragma omp target map (always, to: c, d)
+ #pragma omp teams num_teams (c + 1) thread_limit (d - 1)
+ bar (13);
+ #pragma omp target data map (to: c, d)
+ {
+ #pragma omp target defaultmap (tofrom: scalar)
+ bar2 (14, &c, &d);
+ /* This is one of the cases which can't be generally optimized,
+ the c and d are (or could be) already mapped and whether
+ their device and original values match is unclear. */
+ #pragma omp target map (to: c, d)
+ #pragma omp teams num_teams (c + 1) thread_limit (d - 1)
+ bar (15);
+ }
+ /* This can't be optimized, there are function calls inside of
+ target involved. */
+ #pragma omp target
+ #pragma omp teams num_teams (baz () + 1) thread_limit (baz () - 1)
+ bar (16);
+ #pragma omp target teams num_teams (baz () + 1) thread_limit (baz () - 1)
+ bar (17);
+ /* This one can't be optimized, as v might have different value between
+ host and target. */
+ #pragma omp target
+ #pragma omp teams num_teams (v + 1) thread_limit (v - 1)
+ bar (18);
+}
+
+/* { dg-final { scan-tree-dump-times "num_teams\\(-1\\)" 3 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "thread_limit\\(-1\\)" 3 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "num_teams\\(0\\)" 4 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "thread_limit\\(0\\)" 6 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "num_teams\\(1\\)" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "thread_limit\\(1\\)" 0 "gimple" } } */
--- gcc/testsuite/g++.dg/gomp/target-teams-1.C.jj 2015-10-23 18:45:01.675743377 +0200
+++ gcc/testsuite/g++.dg/gomp/target-teams-1.C 2015-10-23 18:45:49.997049981 +0200
@@ -0,0 +1,92 @@
+// { dg-do compile }
+// { dg-options "-fopenmp -fdump-tree-gimple" }
+
+int v = 6;
+void bar (int);
+void bar2 (int, long *, long *);
+int baz (void);
+#pragma omp declare target to (bar, baz, v)
+
+template <int N>
+void
+foo (int a, int b, long c, long d)
+{
+ /* The OpenMP 4.5 spec says that these expressions are evaluated before
+ target region on combined target teams, so those cases are always
+ fine. */
+ #pragma omp target
+ bar (0);
+ #pragma omp target
+ #pragma omp teams
+ bar (1);
+ #pragma omp target teams
+ bar (2);
+ #pragma omp target
+ #pragma omp teams num_teams (4)
+ bar (3);
+ #pragma omp target teams num_teams (4)
+ bar (4);
+ #pragma omp target
+ #pragma omp teams thread_limit (7)
+ bar (5);
+ #pragma omp target teams thread_limit (7)
+ bar (6);
+ #pragma omp target
+ #pragma omp teams num_teams (4) thread_limit (8)
+ {
+ {
+ bar (7);
+ }
+ }
+ #pragma omp target teams num_teams (4) thread_limit (8)
+ bar (8);
+ #pragma omp target
+ #pragma omp teams num_teams (a) thread_limit (b)
+ bar (9);
+ #pragma omp target teams num_teams (a) thread_limit (b)
+ bar (10);
+ #pragma omp target
+ #pragma omp teams num_teams (c + 1) thread_limit (d - 1)
+ bar (11);
+ #pragma omp target teams num_teams (c + 1) thread_limit (d - 1)
+ bar (12);
+ #pragma omp target map (always, to: c, d)
+ #pragma omp teams num_teams (c + 1) thread_limit (d - 1)
+ bar (13);
+ #pragma omp target data map (to: c, d)
+ {
+ #pragma omp target defaultmap (tofrom: scalar)
+ bar2 (14, &c, &d);
+ /* This is one of the cases which can't be generally optimized,
+ the c and d are (or could be) already mapped and whether
+ their device and original values match is unclear. */
+ #pragma omp target map (to: c, d)
+ #pragma omp teams num_teams (c + 1) thread_limit (d - 1)
+ bar (15);
+ }
+ /* This can't be optimized, there are function calls inside of
+ target involved. */
+ #pragma omp target
+ #pragma omp teams num_teams (baz () + 1) thread_limit (baz () - 1)
+ bar (16);
+ #pragma omp target teams num_teams (baz () + 1) thread_limit (baz () - 1)
+ bar (17);
+ /* This one can't be optimized, as v might have different value between
+ host and target. */
+ #pragma omp target
+ #pragma omp teams num_teams (v + 1) thread_limit (v - 1)
+ bar (18);
+}
+
+void
+foo (int a, int b, long c, long d)
+{
+ foo<0> (a, b, c, d);
+}
+
+/* { dg-final { scan-tree-dump-times "num_teams\\(-1\\)" 3 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "thread_limit\\(-1\\)" 3 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "num_teams\\(0\\)" 4 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "thread_limit\\(0\\)" 6 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "num_teams\\(1\\)" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "thread_limit\\(1\\)" 0 "gimple" } } */
--- include/gomp-constants.h.jj 2015-10-14 10:24:24.000000000 +0200
+++ include/gomp-constants.h 2015-10-26 12:56:26.785551282 +0100
@@ -175,7 +175,7 @@ enum gomp_map_kind
#define GOMP_TASK_FLAG_IF (1 << 10)
#define GOMP_TASK_FLAG_NOGROUP (1 << 11)
-/* GOMP_target{_41,update_41,enter_exit_data} flags argument. */
+/* GOMP_target{_ext,update_ext,enter_exit_data} flags argument. */
#define GOMP_TARGET_FLAG_NOWAIT (1 << 0)
#define GOMP_TARGET_FLAG_EXIT_DATA (1 << 1)
/* Internal to libgomp. */
--- libgomp/target.c.jj 2015-10-16 18:10:49.000000000 +0200
+++ libgomp/target.c 2015-10-26 12:56:05.389864306 +0100
@@ -1209,7 +1209,7 @@ gomp_fini_device (struct gomp_device_des
devicep->is_initialized = false;
}
-/* Host fallback for GOMP_target{,_41} routines. */
+/* Host fallback for GOMP_target{,_ext} routines. */
static void
gomp_target_fallback (void (*fn) (void *), void **hostaddrs)
@@ -1265,7 +1265,7 @@ gomp_target_fallback_firstprivate (void
gomp_target_fallback (fn, hostaddrs);
}
-/* Helper function of GOMP_target{,_41} routines. */
+/* Helper function of GOMP_target{,_ext} routines. */
static void *
gomp_get_target_fn_addr (struct gomp_device_descr *devicep,
@@ -1328,13 +1328,31 @@ GOMP_target (int device, void (*fn) (voi
gomp_unmap_vars (tgt_vars, true);
}
+/* Like GOMP_target, but KINDS is 16-bit, UNUSED is no longer present,
+ and several arguments have been added:
+ FLAGS is a bitmask, see GOMP_TARGET_FLAG_* in gomp-constants.h.
+ DEPEND is array of dependencies, see GOMP_task for details.
+ NUM_TEAMS is positive if GOMP_teams will be called in the body with
+ that value, or 1 if teams construct is not present, or 0, if
+ teams construct does not have num_teams clause and so the choice is
+ implementation defined, and -1 if it can't be determined on the host
+ what value will GOMP_teams have on the device.
+ THREAD_LIMIT similarly is positive if GOMP_teams will be called in the
+ body with that value, or 0, if teams construct does not have thread_limit
+ clause or the teams construct is not present, or -1 if it can't be
+ determined on the host what value will GOMP_teams have on the device. */
+
void
-GOMP_target_41 (int device, void (*fn) (void *), size_t mapnum,
- void **hostaddrs, size_t *sizes, unsigned short *kinds,
- unsigned int flags, void **depend)
+GOMP_target_ext (int device, void (*fn) (void *), size_t mapnum,
+ void **hostaddrs, size_t *sizes, unsigned short *kinds,
+ unsigned int flags, void **depend, int num_teams,
+ int thread_limit)
{
struct gomp_device_descr *devicep = resolve_device (device);
+ (void) num_teams;
+ (void) thread_limit;
+
/* If there are depend clauses, but nowait is not present,
block the parent task until the dependencies are resolved
and then just continue with the rest of the function as if it
@@ -1372,7 +1390,7 @@ GOMP_target_41 (int device, void (*fn) (
gomp_unmap_vars (tgt_vars, true);
}
-/* Host fallback for GOMP_target_data{,_41} routines. */
+/* Host fallback for GOMP_target_data{,_ext} routines. */
static void
gomp_target_data_fallback (void)
@@ -1411,8 +1429,8 @@ GOMP_target_data (int device, const void
}
void
-GOMP_target_data_41 (int device, size_t mapnum, void **hostaddrs, size_t *sizes,
- unsigned short *kinds)
+GOMP_target_data_ext (int device, size_t mapnum, void **hostaddrs,
+ size_t *sizes, unsigned short *kinds)
{
struct gomp_device_descr *devicep = resolve_device (device);
@@ -1454,9 +1472,9 @@ GOMP_target_update (int device, const vo
}
void
-GOMP_target_update_41 (int device, size_t mapnum, void **hostaddrs,
- size_t *sizes, unsigned short *kinds,
- unsigned int flags, void **depend)
+GOMP_target_update_ext (int device, size_t mapnum, void **hostaddrs,
+ size_t *sizes, unsigned short *kinds,
+ unsigned int flags, void **depend)
{
struct gomp_device_descr *devicep = resolve_device (device);
@@ -1648,7 +1666,7 @@ gomp_target_task_fn (void *data)
struct gomp_target_task *ttask = (struct gomp_target_task *) data;
if (ttask->fn != NULL)
{
- /* GOMP_target_41 */
+ /* GOMP_target_ext */
}
else if (ttask->devicep == NULL
|| !(ttask->devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400))
--- libgomp/libgomp.map.jj 2015-10-14 13:41:18.000000000 +0200
+++ libgomp/libgomp.map 2015-10-26 12:35:21.023107290 +0100
@@ -266,9 +266,9 @@ GOMP_4.0.1 {
GOMP_4.5 {
global:
- GOMP_target_41;
- GOMP_target_data_41;
- GOMP_target_update_41;
+ GOMP_target_ext;
+ GOMP_target_data_ext;
+ GOMP_target_update_ext;
GOMP_target_enter_exit_data;
GOMP_taskloop;
GOMP_taskloop_ull;
--- libgomp/libgomp_g.h.jj 2015-10-14 13:53:33.000000000 +0200
+++ libgomp/libgomp_g.h 2015-10-26 12:34:52.825518094 +0100
@@ -277,17 +277,18 @@ extern void GOMP_single_copy_end (void *
extern void GOMP_target (int, void (*) (void *), const void *,
size_t, void **, size_t *, unsigned char *);
-extern void GOMP_target_41 (int, void (*) (void *), size_t, void **, size_t *,
- unsigned short *, unsigned int, void **);
+extern void GOMP_target_ext (int, void (*) (void *), size_t, void **, size_t *,
+ unsigned short *, unsigned int, void **,
+ int, int);
extern void GOMP_target_data (int, const void *,
size_t, void **, size_t *, unsigned char *);
-extern void GOMP_target_data_41 (int, size_t, void **, size_t *,
- unsigned short *);
+extern void GOMP_target_data_ext (int, size_t, void **, size_t *,
+ unsigned short *);
extern void GOMP_target_end_data (void);
extern void GOMP_target_update (int, const void *,
size_t, void **, size_t *, unsigned char *);
-extern void GOMP_target_update_41 (int, size_t, void **, size_t *,
- unsigned short *, unsigned int, void **);
+extern void GOMP_target_update_ext (int, size_t, void **, size_t *,
+ unsigned short *, unsigned int, void **);
extern void GOMP_target_enter_exit_data (int, size_t, void **, size_t *,
unsigned short *, unsigned int,
void **);
--- libgomp/testsuite/libgomp.c/target-teams-1.c.jj 2015-10-26 13:48:11.605281851 +0100
+++ libgomp/testsuite/libgomp.c/target-teams-1.c 2015-10-26 14:46:55.010124006 +0100
@@ -0,0 +1,152 @@
+/* { dg-do run } */
+
+#include <omp.h>
+#include <stdlib.h>
+
+int v = 6;
+
+void
+bar (long *x, long *y)
+{
+ *x += 2;
+ *y += 3;
+}
+
+int
+baz (void)
+{
+ return 5;
+}
+
+#pragma omp declare target to (bar, baz, v)
+
+__attribute__((noinline, noclone)) void
+foo (int a, int b, long c, long d)
+{
+ int err;
+ if (omp_get_num_teams () != 1)
+ abort ();
+ /* The OpenMP 4.5 spec says that these expressions are evaluated before
+ target region on combined target teams, so those cases are always
+ fine. */
+ #pragma omp target map(from: err)
+ err = omp_get_num_teams () != 1;
+ if (err)
+ abort ();
+ #pragma omp target map(from: err)
+ #pragma omp teams
+ err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1;
+ if (err)
+ abort ();
+ #pragma omp target teams map(from: err)
+ err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1;
+ if (err)
+ abort ();
+ #pragma omp target map(from: err)
+ #pragma omp teams num_teams (4)
+ err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
+ || omp_get_num_teams () > 4;
+ if (err)
+ abort ();
+ #pragma omp target teams num_teams (4) map(from: err)
+ err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
+ || omp_get_num_teams () > 4;
+ if (err)
+ abort ();
+ #pragma omp target map(from: err)
+ #pragma omp teams thread_limit (7)
+ err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
+ || omp_get_thread_limit () > 7;
+ if (err)
+ abort ();
+ #pragma omp target teams thread_limit (7) map(from: err)
+ err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
+ || omp_get_thread_limit () > 7;
+ if (err)
+ abort ();
+ #pragma omp target map(from: err)
+ #pragma omp teams num_teams (4) thread_limit (8)
+ {
+ {
+ err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
+ || omp_get_num_teams () > 4 || omp_get_thread_limit () > 8;
+ }
+ }
+ if (err)
+ abort ();
+ #pragma omp target teams num_teams (4) thread_limit (8) map(from: err)
+ err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
+ || omp_get_num_teams () > 4 || omp_get_thread_limit () > 8;
+ if (err)
+ abort ();
+ #pragma omp target map(from: err)
+ #pragma omp teams num_teams (a) thread_limit (b)
+ err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
+ || omp_get_num_teams () > a || omp_get_thread_limit () > b;
+ if (err)
+ abort ();
+ #pragma omp target teams num_teams (a) thread_limit (b) map(from: err)
+ err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
+ || omp_get_num_teams () > a || omp_get_thread_limit () > b;
+ if (err)
+ abort ();
+ #pragma omp target map(from: err)
+ #pragma omp teams num_teams (c + 1) thread_limit (d - 1)
+ err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
+ || omp_get_num_teams () > c + 1 || omp_get_thread_limit () > d - 1;
+ if (err)
+ abort ();
+ #pragma omp target teams num_teams (c + 1) thread_limit (d - 1) map(from: err)
+ err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
+ || omp_get_num_teams () > c + 1 || omp_get_thread_limit () > d - 1;
+ if (err)
+ abort ();
+ #pragma omp target map (always, to: c, d) map(from: err)
+ #pragma omp teams num_teams (c + 1) thread_limit (d - 1)
+ err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
+ || omp_get_num_teams () > c + 1 || omp_get_thread_limit () > d - 1;
+ if (err)
+ abort ();
+ #pragma omp target data map (to: c, d)
+ {
+ #pragma omp target defaultmap (tofrom: scalar)
+ bar (&c, &d);
+ /* This is one of the cases which can't be generally optimized,
+ the c and d are (or could be) already mapped and whether
+ their device and original values match is unclear. */
+ #pragma omp target map (to: c, d) map(from: err)
+ #pragma omp teams num_teams (c + 1) thread_limit (d - 1)
+ err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
+ || omp_get_num_teams () > c + 1 || omp_get_thread_limit () > d - 1;
+ if (err)
+ abort ();
+ }
+ /* This can't be optimized, there are function calls inside of
+ target involved. */
+ #pragma omp target map(from: err)
+ #pragma omp teams num_teams (baz () + 1) thread_limit (baz () - 1)
+ err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
+ || omp_get_num_teams () > baz () + 1 || omp_get_thread_limit () > baz () - 1;
+ if (err)
+ abort ();
+ #pragma omp target teams num_teams (baz () + 1) thread_limit (baz () - 1) map(from: err)
+ err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
+ || omp_get_num_teams () > baz () + 1 || omp_get_thread_limit () > baz () - 1;
+ if (err)
+ abort ();
+ /* This one can't be optimized, as v might have different value between
+ host and target. */
+ #pragma omp target map(from: err)
+ #pragma omp teams num_teams (v + 1) thread_limit (v - 1)
+ err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1
+ || omp_get_num_teams () > v + 1 || omp_get_thread_limit () > v - 1;
+ if (err)
+ abort ();
+}
+
+int
+main ()
+{
+ foo (3, 5, 7, 9);
+ return 0;
+}
Jakub