@@ -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,
@@ -654,11 +654,14 @@ enum cgraph_simd_clone_arg_type
/* These are only for integer/pointer arguments passed by value. */
SIMD_CLONE_ARG_TYPE_LINEAR_CONSTANT_STEP,
SIMD_CLONE_ARG_TYPE_LINEAR_VARIABLE_STEP,
- /* These 3 are only for reference type arguments or arguments passed
+ /* These 6 are only for reference type arguments or arguments passed
by reference. */
SIMD_CLONE_ARG_TYPE_LINEAR_REF_CONSTANT_STEP,
+ SIMD_CLONE_ARG_TYPE_LINEAR_REF_VARIABLE_STEP,
SIMD_CLONE_ARG_TYPE_LINEAR_UVAL_CONSTANT_STEP,
+ SIMD_CLONE_ARG_TYPE_LINEAR_UVAL_VARIABLE_STEP,
SIMD_CLONE_ARG_TYPE_LINEAR_VAL_CONSTANT_STEP,
+ SIMD_CLONE_ARG_TYPE_LINEAR_VAL_VARIABLE_STEP,
SIMD_CLONE_ARG_TYPE_MASK
};
@@ -700,7 +703,7 @@ struct GTY(()) cgraph_simd_clone_arg {
/* For arg_type SIMD_CLONE_ARG_TYPE_LINEAR_*CONSTANT_STEP this is
the constant linear step, if arg_type is
- SIMD_CLONE_ARG_TYPE_LINEAR_VARIABLE_STEP, this is index of
+ SIMD_CLONE_ARG_TYPE_LINEAR_*VARIABLE_STEP, this is index of
the uniform argument holding the step, otherwise 0. */
HOST_WIDE_INT linear_step;
@@ -113,6 +113,14 @@ DEF_GOMP_BUILTIN (BUILT_IN_GOMP_LOOP_RUN
"GOMP_loop_runtime_start",
BT_FN_BOOL_LONG_LONG_LONG_LONGPTR_LONGPTR,
ATTR_NOTHROW_LEAF_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_LOOP_NONMONOTONIC_DYNAMIC_START,
+ "GOMP_loop_nonmonotonic_dynamic_start",
+ BT_FN_BOOL_LONG_LONG_LONG_LONG_LONGPTR_LONGPTR,
+ ATTR_NOTHROW_LEAF_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_LOOP_NONMONOTONIC_GUIDED_START,
+ "GOMP_loop_nonmonotonic_guided_start",
+ BT_FN_BOOL_LONG_LONG_LONG_LONG_LONGPTR_LONGPTR,
+ ATTR_NOTHROW_LEAF_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_GOMP_LOOP_ORDERED_STATIC_START,
"GOMP_loop_ordered_static_start",
BT_FN_BOOL_LONG_LONG_LONG_LONG_LONGPTR_LONGPTR,
@@ -153,6 +161,12 @@ DEF_GOMP_BUILTIN (BUILT_IN_GOMP_LOOP_GUI
BT_FN_BOOL_LONGPTR_LONGPTR, ATTR_NOTHROW_LEAF_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_GOMP_LOOP_RUNTIME_NEXT, "GOMP_loop_runtime_next",
BT_FN_BOOL_LONGPTR_LONGPTR, ATTR_NOTHROW_LEAF_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_LOOP_NONMONOTONIC_DYNAMIC_NEXT,
+ "GOMP_loop_nonmonotonic_dynamic_next",
+ BT_FN_BOOL_LONGPTR_LONGPTR, ATTR_NOTHROW_LEAF_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_LOOP_NONMONOTONIC_GUIDED_NEXT,
+ "GOMP_loop_nonmonotonic_guided_next",
+ BT_FN_BOOL_LONGPTR_LONGPTR, ATTR_NOTHROW_LEAF_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_GOMP_LOOP_ORDERED_STATIC_NEXT,
"GOMP_loop_ordered_static_next",
BT_FN_BOOL_LONGPTR_LONGPTR, ATTR_NOTHROW_LEAF_LIST)
@@ -181,6 +195,14 @@ DEF_GOMP_BUILTIN (BUILT_IN_GOMP_LOOP_ULL
"GOMP_loop_ull_runtime_start",
BT_FN_BOOL_BOOL_ULL_ULL_ULL_ULLPTR_ULLPTR,
ATTR_NOTHROW_LEAF_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_LOOP_ULL_NONMONOTONIC_DYNAMIC_START,
+ "GOMP_loop_ull_nonmonotonic_dynamic_start",
+ BT_FN_BOOL_BOOL_ULL_ULL_ULL_ULL_ULLPTR_ULLPTR,
+ ATTR_NOTHROW_LEAF_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_LOOP_ULL_NONMONOTONIC_GUIDED_START,
+ "GOMP_loop_ull_nonmonotonic_guided_start",
+ BT_FN_BOOL_BOOL_ULL_ULL_ULL_ULL_ULLPTR_ULLPTR,
+ ATTR_NOTHROW_LEAF_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_GOMP_LOOP_ULL_ORDERED_STATIC_START,
"GOMP_loop_ull_ordered_static_start",
BT_FN_BOOL_BOOL_ULL_ULL_ULL_ULL_ULLPTR_ULLPTR,
@@ -213,13 +235,23 @@ DEF_GOMP_BUILTIN (BUILT_IN_GOMP_LOOP_ULL
"GOMP_loop_ull_doacross_runtime_start",
BT_FN_BOOL_UINT_ULLPTR_ULLPTR_ULLPTR,
ATTR_NOTHROW_LEAF_LIST)
-DEF_GOMP_BUILTIN (BUILT_IN_GOMP_LOOP_ULL_STATIC_NEXT, "GOMP_loop_ull_static_next",
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_LOOP_ULL_STATIC_NEXT,
+ "GOMP_loop_ull_static_next",
+ BT_FN_BOOL_ULONGLONGPTR_ULONGLONGPTR, ATTR_NOTHROW_LEAF_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_LOOP_ULL_DYNAMIC_NEXT,
+ "GOMP_loop_ull_dynamic_next",
+ BT_FN_BOOL_ULONGLONGPTR_ULONGLONGPTR, ATTR_NOTHROW_LEAF_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_LOOP_ULL_GUIDED_NEXT,
+ "GOMP_loop_ull_guided_next",
BT_FN_BOOL_ULONGLONGPTR_ULONGLONGPTR, ATTR_NOTHROW_LEAF_LIST)
-DEF_GOMP_BUILTIN (BUILT_IN_GOMP_LOOP_ULL_DYNAMIC_NEXT, "GOMP_loop_ull_dynamic_next",
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_LOOP_ULL_RUNTIME_NEXT,
+ "GOMP_loop_ull_runtime_next",
BT_FN_BOOL_ULONGLONGPTR_ULONGLONGPTR, ATTR_NOTHROW_LEAF_LIST)
-DEF_GOMP_BUILTIN (BUILT_IN_GOMP_LOOP_ULL_GUIDED_NEXT, "GOMP_loop_ull_guided_next",
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_LOOP_ULL_NONMONOTONIC_DYNAMIC_NEXT,
+ "GOMP_loop_ull_nonmonotonic_dynamic_next",
BT_FN_BOOL_ULONGLONGPTR_ULONGLONGPTR, ATTR_NOTHROW_LEAF_LIST)
-DEF_GOMP_BUILTIN (BUILT_IN_GOMP_LOOP_ULL_RUNTIME_NEXT, "GOMP_loop_ull_runtime_next",
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_LOOP_ULL_NONMONOTONIC_GUIDED_NEXT,
+ "GOMP_loop_ull_nonmonotonic_guided_next",
BT_FN_BOOL_ULONGLONGPTR_ULONGLONGPTR, ATTR_NOTHROW_LEAF_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_GOMP_LOOP_ULL_ORDERED_STATIC_NEXT,
"GOMP_loop_ull_ordered_static_next",
@@ -252,6 +284,14 @@ DEF_GOMP_BUILTIN (BUILT_IN_GOMP_PARALLEL
"GOMP_parallel_loop_runtime",
BT_FN_VOID_OMPFN_PTR_UINT_LONG_LONG_LONG_UINT,
ATTR_NOTHROW_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_PARALLEL_LOOP_NONMONOTONIC_DYNAMIC,
+ "GOMP_parallel_loop_nonmonotonic_dynamic",
+ BT_FN_VOID_OMPFN_PTR_UINT_LONG_LONG_LONG_LONG_UINT,
+ ATTR_NOTHROW_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_PARALLEL_LOOP_NONMONOTONIC_GUIDED,
+ "GOMP_parallel_loop_nonmonotonic_guided",
+ BT_FN_VOID_OMPFN_PTR_UINT_LONG_LONG_LONG_LONG_UINT,
+ ATTR_NOTHROW_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_GOMP_LOOP_END, "GOMP_loop_end",
BT_FN_VOID, ATTR_NOTHROW_LEAF_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_GOMP_LOOP_END_CANCEL, "GOMP_loop_end_cancel",
@@ -302,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,
@@ -450,7 +450,10 @@ enum omp_clause_schedule_kind {
OMP_CLAUSE_SCHEDULE_AUTO,
OMP_CLAUSE_SCHEDULE_RUNTIME,
OMP_CLAUSE_SCHEDULE_CILKFOR,
- OMP_CLAUSE_SCHEDULE_LAST
+ OMP_CLAUSE_SCHEDULE_MASK = (1 << 3) - 1,
+ OMP_CLAUSE_SCHEDULE_MONOTONIC = (1 << 3),
+ OMP_CLAUSE_SCHEDULE_NONMONOTONIC = (1 << 4),
+ OMP_CLAUSE_SCHEDULE_LAST = 2 * OMP_CLAUSE_SCHEDULE_NONMONOTONIC - 1
};
enum omp_clause_default_kind {
@@ -1109,28 +1109,28 @@ DEFTREECODE (OMP_TASK, "omp_task", tcc_s
DEFTREECODE (OMP_FOR, "omp_for", tcc_statement, 7)
/* OpenMP - #pragma omp simd [clause1 ... clauseN]
- Operands like operands 1-6 of OMP_FOR. */
-DEFTREECODE (OMP_SIMD, "omp_simd", tcc_statement, 6)
+ Operands like for OMP_FOR. */
+DEFTREECODE (OMP_SIMD, "omp_simd", tcc_statement, 7)
/* Cilk Plus - #pragma simd [clause1 ... clauseN]
- Operands like operands 1-6 of OMP_FOR. */
-DEFTREECODE (CILK_SIMD, "cilk_simd", tcc_statement, 6)
+ Operands like for OMP_FOR. */
+DEFTREECODE (CILK_SIMD, "cilk_simd", tcc_statement, 7)
/* Cilk Plus - _Cilk_for (..)
- Operands like operands 1-6 of OMP_FOR. */
-DEFTREECODE (CILK_FOR, "cilk_for", tcc_statement, 6)
+ Operands like for OMP_FOR. */
+DEFTREECODE (CILK_FOR, "cilk_for", tcc_statement, 7)
/* OpenMP - #pragma omp distribute [clause1 ... clauseN]
- Operands like operands 1-6 of OMP_FOR. */
-DEFTREECODE (OMP_DISTRIBUTE, "omp_distribute", tcc_statement, 6)
+ Operands like for OMP_FOR. */
+DEFTREECODE (OMP_DISTRIBUTE, "omp_distribute", tcc_statement, 7)
/* OpenMP - #pragma omp taskloop [clause1 ... clauseN]
- Operands like operands 1-6 of OMP_FOR. */
-DEFTREECODE (OMP_TASKLOOP, "omp_taskloop", tcc_statement, 6)
+ Operands like for OMP_FOR. */
+DEFTREECODE (OMP_TASKLOOP, "omp_taskloop", tcc_statement, 7)
/* OpenMP - #pragma acc loop [clause1 ... clauseN]
- Operands like operands 1-6 of OMP_FOR. */
-DEFTREECODE (OACC_LOOP, "oacc_loop", tcc_statement, 6)
+ Operands like for OMP_FOR. */
+DEFTREECODE (OACC_LOOP, "oacc_loop", tcc_statement, 7)
/* OpenMP - #pragma omp teams [clause1 ... clauseN]
Operand 0: OMP_TEAMS_BODY: Teams body.
@@ -1259,8 +1259,7 @@ extern void protected_set_expr_location
#define OMP_FOR_COND(NODE) TREE_OPERAND (OMP_LOOP_CHECK (NODE), 3)
#define OMP_FOR_INCR(NODE) TREE_OPERAND (OMP_LOOP_CHECK (NODE), 4)
#define OMP_FOR_PRE_BODY(NODE) TREE_OPERAND (OMP_LOOP_CHECK (NODE), 5)
-/* Note that this is only available for OMP_FOR, hence OMP_FOR_CHECK. */
-#define OMP_FOR_ORIG_DECLS(NODE) TREE_OPERAND (OMP_FOR_CHECK (NODE), 6)
+#define OMP_FOR_ORIG_DECLS(NODE) TREE_OPERAND (OMP_LOOP_CHECK (NODE), 6)
#define OMP_SECTIONS_BODY(NODE) TREE_OPERAND (OMP_SECTIONS_CHECK (NODE), 0)
#define OMP_SECTIONS_CLAUSES(NODE) TREE_OPERAND (OMP_SECTIONS_CHECK (NODE), 1)
@@ -1452,10 +1451,6 @@ extern void protected_set_expr_location
OMP_CLAUSE_MAP with GOMP_MAP_POINTER are marked with this flag. */
#define OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION(NODE) \
(OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.public_flag)
-/* Nonzero if the same decl appears both in OMP_CLAUSE_MAP and either
- OMP_CLAUSE_PRIVATE or OMP_CLAUSE_FIRSTPRIVATE. */
-#define OMP_CLAUSE_MAP_PRIVATE(NODE) \
- TREE_PRIVATE (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP))
/* Nonzero if this is a mapped array section, that might need special
treatment if OMP_CLAUSE_SIZE is zero. */
#define OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION(NODE) \
@@ -432,9 +432,24 @@ dump_omp_clause (pretty_printer *pp, tre
case OMP_CLAUSE_SCHEDULE:
pp_string (pp, "schedule(");
+ if (OMP_CLAUSE_SCHEDULE_KIND (clause)
+ & (OMP_CLAUSE_SCHEDULE_MONOTONIC
+ | OMP_CLAUSE_SCHEDULE_NONMONOTONIC))
+ {
+ if (OMP_CLAUSE_SCHEDULE_KIND (clause)
+ & OMP_CLAUSE_SCHEDULE_MONOTONIC)
+ pp_string (pp, "monotonic");
+ else
+ pp_string (pp, "nonmonotonic");
+ if (OMP_CLAUSE_SCHEDULE_SIMD (clause))
+ pp_comma (pp);
+ else
+ pp_colon (pp);
+ }
if (OMP_CLAUSE_SCHEDULE_SIMD (clause))
pp_string (pp, "simd:");
- switch (OMP_CLAUSE_SCHEDULE_KIND (clause))
+
+ switch (OMP_CLAUSE_SCHEDULE_KIND (clause) & OMP_CLAUSE_SCHEDULE_MASK)
{
case OMP_CLAUSE_SCHEDULE_STATIC:
pp_string (pp, "static");
@@ -630,9 +645,15 @@ dump_omp_clause (pretty_printer *pp, tre
case GOMP_MAP_FIRSTPRIVATE_POINTER:
pp_string (pp, "firstprivate");
break;
+ case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
+ pp_string (pp, "firstprivate ref");
+ break;
case GOMP_MAP_STRUCT:
pp_string (pp, "struct");
break;
+ case GOMP_MAP_ALWAYS_POINTER:
+ pp_string (pp, "always_pointer");
+ break;
default:
gcc_unreachable ();
}
@@ -642,16 +663,22 @@ dump_omp_clause (pretty_printer *pp, tre
print_clause_size:
if (OMP_CLAUSE_SIZE (clause))
{
- if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP
- && (OMP_CLAUSE_MAP_KIND (clause) == GOMP_MAP_POINTER
- || OMP_CLAUSE_MAP_KIND (clause)
- == GOMP_MAP_FIRSTPRIVATE_POINTER))
- pp_string (pp, " [pointer assign, bias: ");
- else if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_KIND (clause) == GOMP_MAP_TO_PSET)
- pp_string (pp, " [pointer set, len: ");
- else
- pp_string (pp, " [len: ");
+ switch (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP
+ ? OMP_CLAUSE_MAP_KIND (clause) : GOMP_MAP_TO)
+ {
+ case GOMP_MAP_POINTER:
+ case GOMP_MAP_FIRSTPRIVATE_POINTER:
+ case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
+ case GOMP_MAP_ALWAYS_POINTER:
+ pp_string (pp, " [pointer assign, bias: ");
+ break;
+ case GOMP_MAP_TO_PSET:
+ pp_string (pp, " [pointer set, len: ");
+ break;
+ default:
+ pp_string (pp, " [len: ");
+ break;
+ }
dump_generic_node (pp, OMP_CLAUSE_SIZE (clause),
spc, flags, false);
pp_right_bracket (pp);
@@ -2826,6 +2826,9 @@ vectorizable_simd_clone_call (gimple *st
case SIMD_CLONE_ARG_TYPE_LINEAR_VARIABLE_STEP:
case SIMD_CLONE_ARG_TYPE_LINEAR_VAL_CONSTANT_STEP:
case SIMD_CLONE_ARG_TYPE_LINEAR_UVAL_CONSTANT_STEP:
+ case SIMD_CLONE_ARG_TYPE_LINEAR_REF_VARIABLE_STEP:
+ case SIMD_CLONE_ARG_TYPE_LINEAR_VAL_VARIABLE_STEP:
+ case SIMD_CLONE_ARG_TYPE_LINEAR_UVAL_VARIABLE_STEP:
/* FORNOW */
i = -1;
break;
@@ -3098,6 +3101,9 @@ vectorizable_simd_clone_call (gimple *st
}
break;
case SIMD_CLONE_ARG_TYPE_LINEAR_VARIABLE_STEP:
+ case SIMD_CLONE_ARG_TYPE_LINEAR_REF_VARIABLE_STEP:
+ case SIMD_CLONE_ARG_TYPE_LINEAR_VAL_VARIABLE_STEP:
+ case SIMD_CLONE_ARG_TYPE_LINEAR_UVAL_VARIABLE_STEP:
default:
gcc_unreachable ();
}
@@ -84,6 +84,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)
@@ -5847,9 +5850,10 @@ omp_default_clause (struct gimplify_omp_
{
splay_tree_node n2;
- if ((octx->region_type & (ORT_TARGET_DATA | ORT_TARGET)) != 0)
- continue;
n2 = splay_tree_lookup (octx->variables, (splay_tree_key) decl);
+ if ((octx->region_type & (ORT_TARGET_DATA | ORT_TARGET)) != 0
+ && (n2 == NULL || (n2->value & GOVD_DATA_SHARE_CLASS) == 0))
+ continue;
if (n2 && (n2->value & GOVD_DATA_SHARE_CLASS) != GOVD_SHARED)
{
flags |= GOVD_FIRSTPRIVATE;
@@ -5964,8 +5968,13 @@ omp_notice_variable (struct gimplify_omp
else if (is_scalar)
nflags |= GOVD_FIRSTPRIVATE;
}
+ tree type = TREE_TYPE (decl);
if (nflags == flags
- && !lang_hooks.types.omp_mappable_type (TREE_TYPE (decl)))
+ && gimplify_omp_ctxp->target_firstprivatize_array_bases
+ && lang_hooks.decls.omp_privatize_by_reference (decl))
+ type = TREE_TYPE (type);
+ if (nflags == flags
+ && !lang_hooks.types.omp_mappable_type (type))
{
error ("%qD referenced in target region does not have "
"a mappable type", decl);
@@ -6084,6 +6093,9 @@ omp_is_private (struct gimplify_omp_ctx
else if ((n->value & GOVD_REDUCTION) != 0)
error ("iteration variable %qE should not be reduction",
DECL_NAME (decl));
+ else if (simd == 0 && (n->value & GOVD_LINEAR) != 0)
+ error ("iteration variable %qE should not be linear",
+ DECL_NAME (decl));
else if (simd == 1 && (n->value & GOVD_LASTPRIVATE) != 0)
error ("iteration variable %qE should not be lastprivate",
DECL_NAME (decl));
@@ -6141,10 +6153,12 @@ omp_check_private (struct gimplify_omp_c
return true;
}
- if ((ctx->region_type & (ORT_TARGET | ORT_TARGET_DATA)) != 0)
+ n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
+
+ if ((ctx->region_type & (ORT_TARGET | ORT_TARGET_DATA)) != 0
+ && (n == NULL || (n->value & GOVD_DATA_SHARE_CLASS) == 0))
continue;
- n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
if (n != NULL)
{
if ((n->value & GOVD_LOCAL) != 0
@@ -6175,12 +6189,12 @@ omp_no_lastprivate (struct gimplify_omp_
if (!ctx->combined_loop)
return false;
if (ctx->distribute)
- return true;
+ return lang_GNU_Fortran ();
break;
case ORT_COMBINED_PARALLEL:
break;
case ORT_COMBINED_TEAMS:
- return true;
+ return lang_GNU_Fortran ();
default:
return false;
}
@@ -6215,7 +6229,7 @@ gimplify_scan_omp_clauses (tree *list_p,
struct gimplify_omp_ctx *ctx, *outer_ctx;
tree c;
hash_map<tree, tree> *struct_map_to_clause = NULL;
- tree *orig_list_p = list_p;
+ tree *prev_list_p = NULL;
ctx = new_omp_context (region_type);
outer_ctx = ctx->outer_context;
@@ -6277,16 +6291,25 @@ gimplify_scan_omp_clauses (tree *list_p,
else if (error_operand_p (decl))
goto do_add;
else if (outer_ctx
- && outer_ctx->region_type == ORT_COMBINED_PARALLEL
+ && (outer_ctx->region_type == ORT_COMBINED_PARALLEL
+ || outer_ctx->region_type == ORT_COMBINED_TEAMS)
&& splay_tree_lookup (outer_ctx->variables,
(splay_tree_key) decl) == NULL)
- omp_add_variable (outer_ctx, decl, GOVD_SHARED | GOVD_SEEN);
+ {
+ omp_add_variable (outer_ctx, decl, GOVD_SHARED | GOVD_SEEN);
+ if (outer_ctx->outer_context)
+ omp_notice_variable (outer_ctx->outer_context, decl, true);
+ }
else if (outer_ctx
&& (outer_ctx->region_type & ORT_TASK) != 0
&& outer_ctx->combined_loop
&& splay_tree_lookup (outer_ctx->variables,
(splay_tree_key) decl) == NULL)
- omp_add_variable (outer_ctx, decl, GOVD_LASTPRIVATE | GOVD_SEEN);
+ {
+ omp_add_variable (outer_ctx, decl, GOVD_LASTPRIVATE | GOVD_SEEN);
+ if (outer_ctx->outer_context)
+ omp_notice_variable (outer_ctx->outer_context, decl, true);
+ }
else if (outer_ctx
&& outer_ctx->region_type == ORT_WORKSHARE
&& outer_ctx->combined_loop
@@ -6300,8 +6323,14 @@ gimplify_scan_omp_clauses (tree *list_p,
== ORT_COMBINED_PARALLEL)
&& splay_tree_lookup (outer_ctx->outer_context->variables,
(splay_tree_key) decl) == NULL)
- omp_add_variable (outer_ctx->outer_context, decl,
- GOVD_SHARED | GOVD_SEEN);
+ {
+ struct gimplify_omp_ctx *octx = outer_ctx->outer_context;
+ omp_add_variable (octx, decl, GOVD_SHARED | GOVD_SEEN);
+ if (octx->outer_context)
+ omp_notice_variable (octx->outer_context, decl, true);
+ }
+ else if (outer_ctx->outer_context)
+ omp_notice_variable (outer_ctx->outer_context, decl, true);
}
goto do_add;
case OMP_CLAUSE_REDUCTION:
@@ -6324,6 +6353,23 @@ gimplify_scan_omp_clauses (tree *list_p,
omp_notice_variable (ctx, v, true);
}
decl = TREE_OPERAND (decl, 0);
+ if (TREE_CODE (decl) == POINTER_PLUS_EXPR)
+ {
+ if (gimplify_expr (&TREE_OPERAND (decl, 1), pre_p,
+ NULL, is_gimple_val, fb_rvalue)
+ == GS_ERROR)
+ {
+ remove = true;
+ break;
+ }
+ v = TREE_OPERAND (decl, 1);
+ if (DECL_P (v))
+ {
+ omp_firstprivatize_variable (ctx, v);
+ omp_notice_variable (ctx, v, true);
+ }
+ decl = TREE_OPERAND (decl, 0);
+ }
if (TREE_CODE (decl) == ADDR_EXPR
|| TREE_CODE (decl) == INDIRECT_REF)
decl = TREE_OPERAND (decl, 0);
@@ -6397,9 +6443,7 @@ gimplify_scan_omp_clauses (tree *list_p,
{
if (octx->outer_context
&& (octx->outer_context->region_type
- == ORT_COMBINED_PARALLEL
- || (octx->outer_context->region_type
- == ORT_COMBINED_TEAMS)))
+ == ORT_COMBINED_PARALLEL))
octx = octx->outer_context;
else if (omp_check_private (octx, decl, false))
break;
@@ -6414,8 +6458,15 @@ gimplify_scan_omp_clauses (tree *list_p,
&& octx == outer_ctx)
flags = GOVD_SEEN | GOVD_SHARED;
else if (octx
+ && octx->region_type == ORT_COMBINED_TEAMS)
+ flags = GOVD_SEEN | GOVD_SHARED;
+ else if (octx
&& octx->region_type == ORT_COMBINED_TARGET)
- flags &= ~GOVD_LASTPRIVATE;
+ {
+ flags &= ~GOVD_LASTPRIVATE;
+ if (flags == GOVD_SEEN)
+ break;
+ }
else
break;
splay_tree_node on
@@ -6458,7 +6509,9 @@ gimplify_scan_omp_clauses (tree *list_p,
case OMP_TARGET_DATA:
case OMP_TARGET_ENTER_DATA:
case OMP_TARGET_EXIT_DATA:
- if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+ || (OMP_CLAUSE_MAP_KIND (c)
+ == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
/* For target {,enter ,exit }data only the array slice is
mapped, but not the pointer to it. */
remove = true;
@@ -6477,7 +6530,9 @@ gimplify_scan_omp_clauses (tree *list_p,
remove = true;
break;
}
- else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+ else if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+ || (OMP_CLAUSE_MAP_KIND (c)
+ == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
&& TREE_CODE (OMP_CLAUSE_SIZE (c)) != INTEGER_CST)
{
OMP_CLAUSE_SIZE (c)
@@ -6536,6 +6591,25 @@ gimplify_scan_omp_clauses (tree *list_p,
break;
}
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER)
+ {
+ /* Error recovery. */
+ if (prev_list_p == NULL)
+ {
+ remove = true;
+ break;
+ }
+ if (OMP_CLAUSE_CHAIN (*prev_list_p) != c)
+ {
+ tree ch = OMP_CLAUSE_CHAIN (*prev_list_p);
+ if (ch == NULL_TREE || OMP_CLAUSE_CHAIN (ch) != c)
+ {
+ remove = true;
+ break;
+ }
+ }
+ }
+
tree offset;
HOST_WIDE_INT bitsize, bitpos;
machine_mode mode;
@@ -6555,56 +6629,64 @@ gimplify_scan_omp_clauses (tree *list_p,
splay_tree_node n
= splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
bool ptr = (OMP_CLAUSE_MAP_KIND (c)
- == GOMP_MAP_FIRSTPRIVATE_POINTER);
- if (n == NULL || (n->value & (ptr ? GOVD_PRIVATE
- : GOVD_MAP)) == 0)
+ == GOMP_MAP_ALWAYS_POINTER);
+ if (n == NULL || (n->value & GOVD_MAP) == 0)
{
+ tree l = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (l, GOMP_MAP_STRUCT);
+ OMP_CLAUSE_DECL (l) = decl;
+ OMP_CLAUSE_SIZE (l) = size_int (1);
+ if (struct_map_to_clause == NULL)
+ struct_map_to_clause = new hash_map<tree, tree>;
+ struct_map_to_clause->put (decl, l);
if (ptr)
{
+ enum gomp_map_kind mkind
+ = code == OMP_TARGET_EXIT_DATA
+ ? GOMP_MAP_RELEASE : GOMP_MAP_ALLOC;
tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
- OMP_CLAUSE_PRIVATE);
- OMP_CLAUSE_DECL (c2) = decl;
- OMP_CLAUSE_CHAIN (c2) = *orig_list_p;
- *orig_list_p = c2;
- if (struct_map_to_clause == NULL)
- struct_map_to_clause = new hash_map<tree, tree>;
- tree *osc;
- if (n == NULL || (n->value & GOVD_MAP) == 0)
- osc = NULL;
- else
- osc = struct_map_to_clause->get (decl);
- if (osc == NULL)
- struct_map_to_clause->put (decl,
- tree_cons (NULL_TREE,
- c,
- NULL_TREE));
- else
- *osc = tree_cons (*osc, c, NULL_TREE);
- flags = GOVD_PRIVATE | GOVD_EXPLICIT;
- goto do_add_decl;
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (c2, mkind);
+ OMP_CLAUSE_DECL (c2)
+ = unshare_expr (OMP_CLAUSE_DECL (c));
+ OMP_CLAUSE_CHAIN (c2) = *prev_list_p;
+ OMP_CLAUSE_SIZE (c2)
+ = TYPE_SIZE_UNIT (ptr_type_node);
+ OMP_CLAUSE_CHAIN (l) = c2;
+ if (OMP_CLAUSE_CHAIN (*prev_list_p) != c)
+ {
+ tree c4 = OMP_CLAUSE_CHAIN (*prev_list_p);
+ tree c3
+ = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (c3, mkind);
+ OMP_CLAUSE_DECL (c3)
+ = unshare_expr (OMP_CLAUSE_DECL (c4));
+ OMP_CLAUSE_SIZE (c3)
+ = TYPE_SIZE_UNIT (ptr_type_node);
+ OMP_CLAUSE_CHAIN (c3) = *prev_list_p;
+ OMP_CLAUSE_CHAIN (c2) = c3;
+ }
+ *prev_list_p = l;
+ prev_list_p = NULL;
+ }
+ else
+ {
+ OMP_CLAUSE_CHAIN (l) = c;
+ *list_p = l;
+ list_p = &OMP_CLAUSE_CHAIN (l);
}
- *list_p = build_omp_clause (OMP_CLAUSE_LOCATION (c),
- OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (*list_p, GOMP_MAP_STRUCT);
- OMP_CLAUSE_DECL (*list_p) = decl;
- OMP_CLAUSE_SIZE (*list_p) = size_int (1);
- OMP_CLAUSE_CHAIN (*list_p) = c;
- if (struct_map_to_clause == NULL)
- struct_map_to_clause = new hash_map<tree, tree>;
- struct_map_to_clause->put (decl, *list_p);
- list_p = &OMP_CLAUSE_CHAIN (*list_p);
flags = GOVD_MAP | GOVD_EXPLICIT;
- if (OMP_CLAUSE_MAP_KIND (c) & GOMP_MAP_FLAG_ALWAYS)
+ if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) || ptr)
flags |= GOVD_SEEN;
goto do_add_decl;
}
else
{
tree *osc = struct_map_to_clause->get (decl);
- tree *sc = NULL, *pt = NULL;
- if (!ptr && TREE_CODE (*osc) == TREE_LIST)
- osc = &TREE_PURPOSE (*osc);
- if (OMP_CLAUSE_MAP_KIND (c) & GOMP_MAP_FLAG_ALWAYS)
+ tree *sc = NULL, *scp = NULL;
+ if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) || ptr)
n->value |= GOVD_SEEN;
offset_int o1, o2;
if (offset)
@@ -6613,18 +6695,16 @@ gimplify_scan_omp_clauses (tree *list_p,
o1 = 0;
if (bitpos)
o1 = o1 + bitpos / BITS_PER_UNIT;
- if (ptr)
- pt = osc;
- else
- sc = &OMP_CLAUSE_CHAIN (*osc);
- for (; ptr ? (*pt && (sc = &TREE_VALUE (*pt)))
- : *sc != c;
- ptr ? (pt = &TREE_CHAIN (*pt))
- : (sc = &OMP_CLAUSE_CHAIN (*sc)))
- if (TREE_CODE (OMP_CLAUSE_DECL (*sc)) != COMPONENT_REF
- && (TREE_CODE (OMP_CLAUSE_DECL (*sc))
- != INDIRECT_REF)
- && TREE_CODE (OMP_CLAUSE_DECL (*sc)) != ARRAY_REF)
+ for (sc = &OMP_CLAUSE_CHAIN (*osc);
+ *sc != c; sc = &OMP_CLAUSE_CHAIN (*sc))
+ if (ptr && sc == prev_list_p)
+ break;
+ else if (TREE_CODE (OMP_CLAUSE_DECL (*sc))
+ != COMPONENT_REF
+ && (TREE_CODE (OMP_CLAUSE_DECL (*sc))
+ != INDIRECT_REF)
+ && (TREE_CODE (OMP_CLAUSE_DECL (*sc))
+ != ARRAY_REF))
break;
else
{
@@ -6653,6 +6733,8 @@ gimplify_scan_omp_clauses (tree *list_p,
&volatilep, false);
if (base != decl)
break;
+ if (scp)
+ continue;
gcc_assert (offset == NULL_TREE
|| TREE_CODE (offset) == INTEGER_CST);
tree d1 = OMP_CLAUSE_DECL (*sc);
@@ -6691,19 +6773,68 @@ gimplify_scan_omp_clauses (tree *list_p,
o2 = o2 + bitpos2 / BITS_PER_UNIT;
if (wi::ltu_p (o1, o2)
|| (wi::eq_p (o1, o2) && bitpos < bitpos2))
- break;
+ {
+ if (ptr)
+ scp = sc;
+ else
+ break;
+ }
}
+ if (remove)
+ break;
+ OMP_CLAUSE_SIZE (*osc)
+ = size_binop (PLUS_EXPR, OMP_CLAUSE_SIZE (*osc),
+ size_one_node);
if (ptr)
{
- if (!remove)
- *pt = tree_cons (TREE_PURPOSE (*osc), c, *pt);
- break;
+ tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+ OMP_CLAUSE_MAP);
+ tree cl = NULL_TREE;
+ enum gomp_map_kind mkind
+ = code == OMP_TARGET_EXIT_DATA
+ ? GOMP_MAP_RELEASE : GOMP_MAP_ALLOC;
+ OMP_CLAUSE_SET_MAP_KIND (c2, mkind);
+ OMP_CLAUSE_DECL (c2)
+ = unshare_expr (OMP_CLAUSE_DECL (c));
+ OMP_CLAUSE_CHAIN (c2) = scp ? *scp : *prev_list_p;
+ OMP_CLAUSE_SIZE (c2)
+ = TYPE_SIZE_UNIT (ptr_type_node);
+ cl = scp ? *prev_list_p : c2;
+ if (OMP_CLAUSE_CHAIN (*prev_list_p) != c)
+ {
+ tree c4 = OMP_CLAUSE_CHAIN (*prev_list_p);
+ tree c3
+ = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (c3, mkind);
+ OMP_CLAUSE_DECL (c3)
+ = unshare_expr (OMP_CLAUSE_DECL (c4));
+ OMP_CLAUSE_SIZE (c3)
+ = TYPE_SIZE_UNIT (ptr_type_node);
+ OMP_CLAUSE_CHAIN (c3) = *prev_list_p;
+ if (!scp)
+ OMP_CLAUSE_CHAIN (c2) = c3;
+ else
+ cl = c3;
+ }
+ if (scp)
+ *scp = c2;
+ if (sc == prev_list_p)
+ {
+ *sc = cl;
+ prev_list_p = NULL;
+ }
+ else
+ {
+ *prev_list_p = OMP_CLAUSE_CHAIN (c);
+ list_p = prev_list_p;
+ prev_list_p = NULL;
+ OMP_CLAUSE_CHAIN (c) = *sc;
+ *sc = cl;
+ continue;
+ }
}
- if (!remove)
- OMP_CLAUSE_SIZE (*osc)
- = size_binop (PLUS_EXPR, OMP_CLAUSE_SIZE (*osc),
- size_one_node);
- if (!remove && *sc != c)
+ else if (*sc != c)
{
*list_p = OMP_CLAUSE_CHAIN (c);
OMP_CLAUSE_CHAIN (c) = *sc;
@@ -6712,9 +6843,19 @@ gimplify_scan_omp_clauses (tree *list_p,
}
}
}
+ if (!remove
+ && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER
+ && OMP_CLAUSE_CHAIN (c)
+ && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c)) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
+ == GOMP_MAP_ALWAYS_POINTER))
+ prev_list_p = 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:
@@ -6923,7 +7064,12 @@ gimplify_scan_omp_clauses (tree *list_p,
|| decl == OMP_CLAUSE_DECL (c)
|| (TREE_CODE (OMP_CLAUSE_DECL (c)) == MEM_REF
&& (TREE_CODE (TREE_OPERAND (OMP_CLAUSE_DECL (c), 0))
- == ADDR_EXPR)))
+ == ADDR_EXPR
+ || (TREE_CODE (TREE_OPERAND (OMP_CLAUSE_DECL (c), 0))
+ == POINTER_PLUS_EXPR
+ && (TREE_CODE (TREE_OPERAND (TREE_OPERAND
+ (OMP_CLAUSE_DECL (c), 0), 0))
+ == ADDR_EXPR)))))
&& omp_check_private (ctx, decl, false))
{
error ("%s variable %qE is private in outer context",
@@ -7192,6 +7338,25 @@ gimplify_adjust_omp_clauses_1 (splay_tre
OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (clause);
OMP_CLAUSE_CHAIN (clause) = nc;
}
+ else if (gimplify_omp_ctxp->target_firstprivatize_array_bases
+ && lang_hooks.decls.omp_privatize_by_reference (decl))
+ {
+ OMP_CLAUSE_DECL (clause) = build_simple_mem_ref (decl);
+ OMP_CLAUSE_SIZE (clause)
+ = unshare_expr (TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (decl))));
+ struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
+ gimplify_omp_ctxp = ctx->outer_context;
+ gimplify_expr (&OMP_CLAUSE_SIZE (clause),
+ pre_p, NULL, is_gimple_val, fb_rvalue);
+ gimplify_omp_ctxp = ctx;
+ tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (clause),
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_DECL (nc) = decl;
+ OMP_CLAUSE_SIZE (nc) = size_zero_node;
+ OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_FIRSTPRIVATE_REFERENCE);
+ OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (clause);
+ OMP_CLAUSE_CHAIN (clause) = nc;
+ }
else
OMP_CLAUSE_SIZE (clause) = DECL_SIZE_UNIT (decl);
}
@@ -7265,6 +7430,15 @@ gimplify_adjust_omp_clauses (gimple_seq
else
OMP_CLAUSE_CODE (c) = OMP_CLAUSE_PRIVATE;
}
+ else if (code == OMP_DISTRIBUTE
+ && OMP_CLAUSE_LASTPRIVATE_FIRSTPRIVATE (c))
+ {
+ remove = true;
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "same variable used in %<firstprivate%> and "
+ "%<lastprivate%> clauses on %<distribute%> "
+ "construct");
+ }
break;
case OMP_CLAUSE_ALIGNED:
@@ -7310,6 +7484,12 @@ gimplify_adjust_omp_clauses (gimple_seq
break;
case OMP_CLAUSE_MAP:
+ if (code == OMP_TARGET_EXIT_DATA
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER)
+ {
+ remove = true;
+ break;
+ }
decl = OMP_CLAUSE_DECL (c);
if (!DECL_P (decl))
{
@@ -7339,8 +7519,7 @@ gimplify_adjust_omp_clauses (gimple_seq
n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
if ((ctx->region_type & ORT_TARGET) != 0
&& !(n->value & GOVD_SEEN)
- && ((OMP_CLAUSE_MAP_KIND (c) & GOMP_MAP_FLAG_ALWAYS) == 0
- || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT))
+ && GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) == 0)
{
remove = true;
/* For struct element mapping, if struct is never referenced
@@ -7361,7 +7540,9 @@ gimplify_adjust_omp_clauses (gimple_seq
else if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_POINTER
- && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER)
+ && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER
+ && (OMP_CLAUSE_MAP_KIND (c)
+ != GOMP_MAP_FIRSTPRIVATE_REFERENCE))
{
/* For GOMP_MAP_FORCE_DEVICEPTR, we'll never enter here, because
for these, TREE_CODE (DECL_SIZE (decl)) will always be
@@ -7404,9 +7585,9 @@ gimplify_adjust_omp_clauses (gimple_seq
{
if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
OMP_CLAUSE_SIZE (c) = DECL_SIZE_UNIT (decl);
- if ((n->value & GOVD_SEEN)
- && (n->value & (GOVD_PRIVATE | GOVD_FIRSTPRIVATE)))
- OMP_CLAUSE_MAP_PRIVATE (c) = 1;
+ gcc_assert ((n->value & GOVD_SEEN) == 0
+ || ((n->value & (GOVD_PRIVATE | GOVD_FIRSTPRIVATE))
+ == 0));
}
break;
@@ -7878,6 +8059,26 @@ gimplify_omp_for (tree *expr_p, gimple_s
OMP_CLAUSE_LINEAR_NO_COPYOUT (c) = 1;
flags |= GOVD_LINEAR_LASTPRIVATE_NO_OUTER;
}
+ else
+ {
+ struct gimplify_omp_ctx *octx = outer->outer_context;
+ if (octx
+ && octx->region_type == ORT_COMBINED_PARALLEL
+ && octx->outer_context
+ && (octx->outer_context->region_type
+ == ORT_WORKSHARE)
+ && octx->outer_context->combined_loop)
+ {
+ octx = octx->outer_context;
+ n = splay_tree_lookup (octx->variables,
+ (splay_tree_key)decl);
+ if (n != NULL && (n->value & GOVD_LOCAL) != 0)
+ {
+ OMP_CLAUSE_LINEAR_NO_COPYOUT (c) = 1;
+ flags |= GOVD_LINEAR_LASTPRIVATE_NO_OUTER;
+ }
+ }
+ }
}
}
@@ -7912,7 +8113,41 @@ gimplify_omp_for (tree *expr_p, gimple_s
{
omp_add_variable (outer, decl,
GOVD_LASTPRIVATE | GOVD_SEEN);
- if (outer->outer_context)
+ if (outer->region_type == ORT_COMBINED_PARALLEL
+ && outer->outer_context
+ && (outer->outer_context->region_type
+ == ORT_WORKSHARE)
+ && outer->outer_context->combined_loop)
+ {
+ outer = outer->outer_context;
+ n = splay_tree_lookup (outer->variables,
+ (splay_tree_key)decl);
+ if (omp_check_private (outer, decl, false))
+ outer = NULL;
+ else if (n == NULL
+ || ((n->value & GOVD_DATA_SHARE_CLASS)
+ == 0))
+ omp_add_variable (outer, decl,
+ GOVD_LASTPRIVATE
+ | GOVD_SEEN);
+ else
+ outer = NULL;
+ }
+ if (outer && outer->outer_context
+ && (outer->outer_context->region_type
+ == ORT_COMBINED_TEAMS))
+ {
+ outer = outer->outer_context;
+ n = splay_tree_lookup (outer->variables,
+ (splay_tree_key)decl);
+ if (n == NULL
+ || (n->value & GOVD_DATA_SHARE_CLASS) == 0)
+ omp_add_variable (outer, decl,
+ GOVD_SHARED | GOVD_SEEN);
+ else
+ outer = NULL;
+ }
+ if (outer && outer->outer_context)
omp_notice_variable (outer->outer_context, decl,
true);
}
@@ -7961,7 +8196,41 @@ gimplify_omp_for (tree *expr_p, gimple_s
{
omp_add_variable (outer, decl,
GOVD_LASTPRIVATE | GOVD_SEEN);
- if (outer->outer_context)
+ if (outer->region_type == ORT_COMBINED_PARALLEL
+ && outer->outer_context
+ && (outer->outer_context->region_type
+ == ORT_WORKSHARE)
+ && outer->outer_context->combined_loop)
+ {
+ outer = outer->outer_context;
+ n = splay_tree_lookup (outer->variables,
+ (splay_tree_key)decl);
+ if (omp_check_private (outer, decl, false))
+ outer = NULL;
+ else if (n == NULL
+ || ((n->value & GOVD_DATA_SHARE_CLASS)
+ == 0))
+ omp_add_variable (outer, decl,
+ GOVD_LASTPRIVATE
+ | GOVD_SEEN);
+ else
+ outer = NULL;
+ }
+ if (outer && outer->outer_context
+ && (outer->outer_context->region_type
+ == ORT_COMBINED_TEAMS))
+ {
+ outer = outer->outer_context;
+ n = splay_tree_lookup (outer->variables,
+ (splay_tree_key)decl);
+ if (n == NULL
+ || (n->value & GOVD_DATA_SHARE_CLASS) == 0)
+ omp_add_variable (outer, decl,
+ GOVD_SHARED | GOVD_SEEN);
+ else
+ outer = NULL;
+ }
+ if (outer && outer->outer_context)
omp_notice_variable (outer->outer_context, decl,
true);
}
@@ -8399,6 +8668,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
@@ -8434,6 +8898,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 ();
@@ -127,6 +127,9 @@ struct omp_region
/* Schedule kind, only used for GIMPLE_OMP_FOR type regions. */
enum omp_clause_schedule_kind sched_kind;
+ /* Schedule modifiers. */
+ unsigned char sched_modifiers;
+
/* True if this is a combined parallel+workshare region. */
bool is_combined_parallel;
@@ -209,6 +212,7 @@ struct omp_for_data
int collapse;
int ordered;
bool have_nowait, have_ordered, simd_schedule;
+ unsigned char sched_modifiers;
enum omp_clause_schedule_kind sched_kind;
struct omp_for_data_loop *loops;
};
@@ -263,6 +267,7 @@ static int target_nesting_level;
static struct omp_region *root_omp_region;
static bitmap task_shared_vars;
static vec<omp_context *> taskreg_contexts;
+static bool omp_any_child_fn_dumped;
static void scan_omp (gimple_seq *, omp_context *);
static tree scan_omp_1_op (tree *, int *, void *);
@@ -478,6 +483,7 @@ extract_omp_for_data (gomp_for *for_stmt
fd->collapse = 1;
fd->ordered = 0;
fd->sched_kind = OMP_CLAUSE_SCHEDULE_STATIC;
+ fd->sched_modifiers = 0;
fd->chunk_size = NULL_TREE;
fd->simd_schedule = false;
if (gimple_omp_for_kind (fd->for_stmt) == GF_OMP_FOR_KIND_CILKFOR)
@@ -498,7 +504,11 @@ extract_omp_for_data (gomp_for *for_stmt
break;
case OMP_CLAUSE_SCHEDULE:
gcc_assert (!distribute && !taskloop);
- fd->sched_kind = OMP_CLAUSE_SCHEDULE_KIND (t);
+ fd->sched_kind
+ = (enum omp_clause_schedule_kind)
+ (OMP_CLAUSE_SCHEDULE_KIND (t) & OMP_CLAUSE_SCHEDULE_MASK);
+ fd->sched_modifiers = (OMP_CLAUSE_SCHEDULE_KIND (t)
+ & ~OMP_CLAUSE_SCHEDULE_MASK);
fd->chunk_size = OMP_CLAUSE_SCHEDULE_CHUNK_EXPR (t);
fd->simd_schedule = OMP_CLAUSE_SCHEDULE_SIMD (t);
break;
@@ -987,7 +997,8 @@ determine_parallel_type (struct omp_regi
tree clauses = gimple_omp_for_clauses (ws_stmt);
tree c = find_omp_clause (clauses, OMP_CLAUSE_SCHEDULE);
if (c == NULL
- || OMP_CLAUSE_SCHEDULE_KIND (c) == OMP_CLAUSE_SCHEDULE_STATIC
+ || ((OMP_CLAUSE_SCHEDULE_KIND (c) & OMP_CLAUSE_SCHEDULE_MASK)
+ == OMP_CLAUSE_SCHEDULE_STATIC)
|| find_omp_clause (clauses, OMP_CLAUSE_ORDERED))
{
region->is_combined_parallel = false;
@@ -1860,6 +1871,8 @@ scan_sharing_clauses (tree clauses, omp_
&& TREE_CODE (decl) == MEM_REF)
{
tree t = TREE_OPERAND (decl, 0);
+ if (TREE_CODE (t) == POINTER_PLUS_EXPR)
+ t = TREE_OPERAND (t, 0);
if (TREE_CODE (t) == INDIRECT_REF
|| TREE_CODE (t) == ADDR_EXPR)
t = TREE_OPERAND (t, 0);
@@ -1999,7 +2012,9 @@ scan_sharing_clauses (tree clauses, omp_
directly. */
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& DECL_P (decl)
- && (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER
+ && ((OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER
+ && (OMP_CLAUSE_MAP_KIND (c)
+ != GOMP_MAP_FIRSTPRIVATE_REFERENCE))
|| TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
&& is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
&& varpool_node::get_create (decl)->offloadable)
@@ -2015,7 +2030,9 @@ scan_sharing_clauses (tree clauses, omp_
break;
}
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+ || (OMP_CLAUSE_MAP_KIND (c)
+ == GOMP_MAP_FIRSTPRIVATE_REFERENCE)))
{
if (TREE_CODE (decl) == COMPONENT_REF
|| (TREE_CODE (decl) == INDIRECT_REF
@@ -2044,11 +2061,7 @@ scan_sharing_clauses (tree clauses, omp_
gcc_assert (TREE_CODE (decl2) == INDIRECT_REF);
decl2 = TREE_OPERAND (decl2, 0);
gcc_assert (DECL_P (decl2));
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_PRIVATE (c))
- install_var_field (decl2, true, 11, ctx);
- else
- install_var_field (decl2, true, 3, ctx);
+ install_var_field (decl2, true, 3, ctx);
install_var_local (decl2, ctx);
install_var_local (decl, ctx);
}
@@ -2059,9 +2072,6 @@ scan_sharing_clauses (tree clauses, omp_
&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
&& TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
install_var_field (decl, true, 7, ctx);
- else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_PRIVATE (c))
- install_var_field (decl, true, 11, ctx);
else
install_var_field (decl, true, 3, ctx);
if (is_gimple_omp_offloaded (ctx->stmt))
@@ -2225,7 +2235,9 @@ scan_sharing_clauses (tree clauses, omp_
break;
decl = OMP_CLAUSE_DECL (c);
if (DECL_P (decl)
- && (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER
+ && ((OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER
+ && (OMP_CLAUSE_MAP_KIND (c)
+ != GOMP_MAP_FIRSTPRIVATE_REFERENCE))
|| TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
&& is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
&& varpool_node::get_create (decl)->offloadable)
@@ -2562,12 +2574,15 @@ add_taskreg_looptemp_clauses (enum gf_ma
&& TREE_CODE (fd.loop.n2) != INTEGER_CST)
{
count += fd.collapse - 1;
- /* For taskloop, if there are lastprivate clauses on the inner
+ /* If there are lastprivate clauses on the inner
GIMPLE_OMP_FOR, add one more temporaries for the total number
of iterations (product of count1 ... countN-1). */
- if (msk == GF_OMP_FOR_KIND_TASKLOOP
- && find_omp_clause (gimple_omp_for_clauses (for_stmt),
- OMP_CLAUSE_LASTPRIVATE))
+ if (find_omp_clause (gimple_omp_for_clauses (for_stmt),
+ OMP_CLAUSE_LASTPRIVATE))
+ count++;
+ else if (msk == GF_OMP_FOR_KIND_FOR
+ && find_omp_clause (gimple_omp_parallel_clauses (stmt),
+ OMP_CLAUSE_LASTPRIVATE))
count++;
}
for (i = 0; i < count; i++)
@@ -4254,6 +4269,8 @@ lower_rec_input_clauses (tree clauses, g
if (c_kind == OMP_CLAUSE_REDUCTION && TREE_CODE (var) == MEM_REF)
{
var = TREE_OPERAND (var, 0);
+ if (TREE_CODE (var) == POINTER_PLUS_EXPR)
+ var = TREE_OPERAND (var, 0);
if (TREE_CODE (var) == INDIRECT_REF
|| TREE_CODE (var) == ADDR_EXPR)
var = TREE_OPERAND (var, 0);
@@ -4282,7 +4299,28 @@ lower_rec_input_clauses (tree clauses, g
if (pass == 0)
continue;
+ tree bias = TREE_OPERAND (OMP_CLAUSE_DECL (c), 1);
tree orig_var = TREE_OPERAND (OMP_CLAUSE_DECL (c), 0);
+ if (TREE_CODE (orig_var) == POINTER_PLUS_EXPR)
+ {
+ tree b = TREE_OPERAND (orig_var, 1);
+ b = maybe_lookup_decl (b, ctx);
+ if (b == NULL)
+ {
+ b = TREE_OPERAND (orig_var, 1);
+ b = maybe_lookup_decl_in_outer_ctx (b, ctx);
+ }
+ if (integer_zerop (bias))
+ bias = b;
+ else
+ {
+ bias = fold_convert_loc (clause_loc,
+ TREE_TYPE (b), bias);
+ bias = fold_build2_loc (clause_loc, PLUS_EXPR,
+ TREE_TYPE (b), b, bias);
+ }
+ orig_var = TREE_OPERAND (orig_var, 0);
+ }
if (TREE_CODE (orig_var) == INDIRECT_REF
|| TREE_CODE (orig_var) == ADDR_EXPR)
orig_var = TREE_OPERAND (orig_var, 0);
@@ -4323,7 +4361,24 @@ lower_rec_input_clauses (tree clauses, g
tree y = create_tmp_var (ptype, name);
gimplify_assign (y, x, ilist);
x = y;
- if (TREE_CODE (TREE_OPERAND (d, 0)) == ADDR_EXPR)
+ tree yb = y;
+
+ if (!integer_zerop (bias))
+ {
+ bias = fold_convert_loc (clause_loc, sizetype, bias);
+ bias = fold_build1_loc (clause_loc, NEGATE_EXPR,
+ sizetype, bias);
+ x = fold_build2_loc (clause_loc, POINTER_PLUS_EXPR,
+ TREE_TYPE (x), x, bias);
+ yb = create_tmp_var (ptype, name);
+ gimplify_assign (yb, x, ilist);
+ x = yb;
+ }
+
+ d = TREE_OPERAND (d, 0);
+ if (TREE_CODE (d) == POINTER_PLUS_EXPR)
+ d = TREE_OPERAND (d, 0);
+ if (TREE_CODE (d) == ADDR_EXPR)
{
if (orig_var != var)
{
@@ -4349,11 +4404,11 @@ lower_rec_input_clauses (tree clauses, g
else
{
gcc_assert (orig_var == var);
- if (TREE_CODE (TREE_OPERAND (d, 0)) == INDIRECT_REF)
+ if (TREE_CODE (d) == INDIRECT_REF)
{
x = create_tmp_var (ptype, name);
TREE_ADDRESSABLE (x) = 1;
- gimplify_assign (x, y, ilist);
+ gimplify_assign (x, yb, ilist);
x = build_fold_addr_expr_loc (clause_loc, x);
}
x = fold_convert_loc (clause_loc, TREE_TYPE (new_var), x);
@@ -4370,9 +4425,9 @@ lower_rec_input_clauses (tree clauses, g
gimplify_assign (y2, y, ilist);
tree ref = build_outer_var_ref (var, ctx);
/* For ref build_outer_var_ref already performs this. */
- if (TREE_CODE (TREE_OPERAND (d, 0)) == INDIRECT_REF)
+ if (TREE_CODE (d) == INDIRECT_REF)
gcc_assert (is_reference (var));
- else if (TREE_CODE (TREE_OPERAND (d, 0)) == ADDR_EXPR)
+ else if (TREE_CODE (d) == ADDR_EXPR)
ref = build_fold_addr_expr (ref);
else if (is_reference (var))
ref = build_fold_addr_expr (ref);
@@ -5467,6 +5522,8 @@ lower_reduction_clauses (tree clauses, g
if (TREE_CODE (var) == MEM_REF)
{
var = TREE_OPERAND (var, 0);
+ if (TREE_CODE (var) == POINTER_PLUS_EXPR)
+ var = TREE_OPERAND (var, 0);
if (TREE_CODE (var) == INDIRECT_REF
|| TREE_CODE (var) == ADDR_EXPR)
var = TREE_OPERAND (var, 0);
@@ -5509,14 +5566,35 @@ lower_reduction_clauses (tree clauses, g
tree v = TYPE_MAX_VALUE (TYPE_DOMAIN (type));
tree i = create_tmp_var (TREE_TYPE (v), NULL);
tree ptype = build_pointer_type (TREE_TYPE (type));
+ tree bias = TREE_OPERAND (d, 1);
+ d = TREE_OPERAND (d, 0);
+ if (TREE_CODE (d) == POINTER_PLUS_EXPR)
+ {
+ tree b = TREE_OPERAND (d, 1);
+ b = maybe_lookup_decl (b, ctx);
+ if (b == NULL)
+ {
+ b = TREE_OPERAND (d, 1);
+ b = maybe_lookup_decl_in_outer_ctx (b, ctx);
+ }
+ if (integer_zerop (bias))
+ bias = b;
+ else
+ {
+ bias = fold_convert_loc (clause_loc, TREE_TYPE (b), bias);
+ bias = fold_build2_loc (clause_loc, PLUS_EXPR,
+ TREE_TYPE (b), b, bias);
+ }
+ d = TREE_OPERAND (d, 0);
+ }
/* For ref build_outer_var_ref already performs this, so
only new_var needs a dereference. */
- if (TREE_CODE (TREE_OPERAND (d, 0)) == INDIRECT_REF)
+ if (TREE_CODE (d) == INDIRECT_REF)
{
new_var = build_simple_mem_ref_loc (clause_loc, new_var);
gcc_assert (is_reference (var) && var == orig_var);
}
- else if (TREE_CODE (TREE_OPERAND (d, 0)) == ADDR_EXPR)
+ else if (TREE_CODE (d) == ADDR_EXPR)
{
if (orig_var == var)
{
@@ -5539,6 +5617,15 @@ lower_reduction_clauses (tree clauses, g
v = maybe_lookup_decl_in_outer_ctx (v, ctx);
gimplify_expr (&v, stmt_seqp, NULL, is_gimple_val, fb_rvalue);
}
+ if (!integer_zerop (bias))
+ {
+ bias = fold_convert_loc (clause_loc, sizetype, bias);
+ new_var = fold_build2_loc (clause_loc, POINTER_PLUS_EXPR,
+ TREE_TYPE (new_var), new_var,
+ unshare_expr (bias));
+ ref = fold_build2_loc (clause_loc, POINTER_PLUS_EXPR,
+ TREE_TYPE (ref), ref, bias);
+ }
new_var = fold_convert_loc (clause_loc, ptype, new_var);
ref = fold_convert_loc (clause_loc, ptype, ref);
tree m = create_tmp_var (ptype, NULL);
@@ -5731,6 +5818,8 @@ lower_send_clauses (tree clauses, gimple
&& TREE_CODE (val) == MEM_REF)
{
val = TREE_OPERAND (val, 0);
+ if (TREE_CODE (val) == POINTER_PLUS_EXPR)
+ val = TREE_OPERAND (val, 0);
if (TREE_CODE (val) == INDIRECT_REF
|| TREE_CODE (val) == ADDR_EXPR)
val = TREE_OPERAND (val, 0);
@@ -6120,11 +6209,26 @@ expand_parallel_call (struct omp_region
{
case GIMPLE_OMP_FOR:
gcc_assert (region->inner->sched_kind != OMP_CLAUSE_SCHEDULE_AUTO);
- start_ix2 = ((int)BUILT_IN_GOMP_PARALLEL_LOOP_STATIC
- + (region->inner->sched_kind
- == OMP_CLAUSE_SCHEDULE_RUNTIME
- ? 3 : region->inner->sched_kind));
- start_ix = (enum built_in_function)start_ix2;
+ switch (region->inner->sched_kind)
+ {
+ case OMP_CLAUSE_SCHEDULE_RUNTIME:
+ start_ix2 = 3;
+ break;
+ case OMP_CLAUSE_SCHEDULE_DYNAMIC:
+ case OMP_CLAUSE_SCHEDULE_GUIDED:
+ if (region->inner->sched_modifiers
+ & OMP_CLAUSE_SCHEDULE_NONMONOTONIC)
+ {
+ start_ix2 = 3 + region->inner->sched_kind;
+ break;
+ }
+ /* FALLTHRU */
+ default:
+ start_ix2 = region->inner->sched_kind;
+ break;
+ }
+ start_ix2 += (int) BUILT_IN_GOMP_PARALLEL_LOOP_STATIC;
+ start_ix = (enum built_in_function) start_ix2;
break;
case GIMPLE_OMP_SECTIONS:
start_ix = BUILT_IN_GOMP_PARALLEL_SECTIONS;
@@ -6941,9 +7045,15 @@ expand_omp_taskreg (struct omp_region *r
node->parallelized_function = 1;
cgraph_node::add_new_function (child_fn, true);
+ bool need_asm = DECL_ASSEMBLER_NAME_SET_P (current_function_decl)
+ && !DECL_ASSEMBLER_NAME_SET_P (child_fn);
+
/* Fix the callgraph edges for child_cfun. Those for cfun will be
fixed in a following pass. */
push_cfun (child_cfun);
+ if (need_asm)
+ assign_assembler_name_if_neeeded (child_fn);
+
if (optimize)
optimize_omp_library_calls (entry_stmt);
cgraph_edge::rebuild_edges ();
@@ -6967,6 +7077,13 @@ expand_omp_taskreg (struct omp_region *r
if (flag_checking && !loops_state_satisfies_p (LOOPS_NEED_FIXUP))
verify_loop_structure ();
pop_cfun ();
+
+ if (dump_file && !gimple_in_ssa_p (cfun))
+ {
+ omp_any_child_fn_dumped = true;
+ dump_function_header (dump_file, child_fn, dump_flags);
+ dump_function_to_file (child_fn, dump_file, dump_flags);
+ }
}
/* Emit a library call to launch the children threads. */
@@ -8968,6 +9085,30 @@ expand_omp_for_static_nochunk (struct om
OMP_CLAUSE__LOOPTEMP_);
gcc_assert (innerc);
endvar = OMP_CLAUSE_DECL (innerc);
+ if (fd->collapse > 1 && TREE_CODE (fd->loop.n2) != INTEGER_CST
+ && gimple_omp_for_kind (fd->for_stmt) == GF_OMP_FOR_KIND_DISTRIBUTE)
+ {
+ int i;
+ for (i = 1; i < fd->collapse; i++)
+ {
+ innerc = find_omp_clause (OMP_CLAUSE_CHAIN (innerc),
+ OMP_CLAUSE__LOOPTEMP_);
+ gcc_assert (innerc);
+ }
+ innerc = find_omp_clause (OMP_CLAUSE_CHAIN (innerc),
+ OMP_CLAUSE__LOOPTEMP_);
+ if (innerc)
+ {
+ /* If needed (distribute parallel for with lastprivate),
+ propagate down the total number of iterations. */
+ tree t = fold_convert (TREE_TYPE (OMP_CLAUSE_DECL (innerc)),
+ fd->loop.n2);
+ t = force_gimple_operand_gsi (&gsi, t, false, NULL_TREE, false,
+ GSI_CONTINUE_LINKING);
+ assign_stmt = gimple_build_assign (OMP_CLAUSE_DECL (innerc), t);
+ gsi_insert_after (&gsi, assign_stmt, GSI_CONTINUE_LINKING);
+ }
+ }
}
t = fold_convert (itype, s0);
t = fold_build2 (MULT_EXPR, itype, t, step);
@@ -9442,6 +9583,30 @@ expand_omp_for_static_chunk (struct omp_
OMP_CLAUSE__LOOPTEMP_);
gcc_assert (innerc);
endvar = OMP_CLAUSE_DECL (innerc);
+ if (fd->collapse > 1 && TREE_CODE (fd->loop.n2) != INTEGER_CST
+ && gimple_omp_for_kind (fd->for_stmt) == GF_OMP_FOR_KIND_DISTRIBUTE)
+ {
+ int i;
+ for (i = 1; i < fd->collapse; i++)
+ {
+ innerc = find_omp_clause (OMP_CLAUSE_CHAIN (innerc),
+ OMP_CLAUSE__LOOPTEMP_);
+ gcc_assert (innerc);
+ }
+ innerc = find_omp_clause (OMP_CLAUSE_CHAIN (innerc),
+ OMP_CLAUSE__LOOPTEMP_);
+ if (innerc)
+ {
+ /* If needed (distribute parallel for with lastprivate),
+ propagate down the total number of iterations. */
+ tree t = fold_convert (TREE_TYPE (OMP_CLAUSE_DECL (innerc)),
+ fd->loop.n2);
+ t = force_gimple_operand_gsi (&gsi, t, false, NULL_TREE, false,
+ GSI_CONTINUE_LINKING);
+ assign_stmt = gimple_build_assign (OMP_CLAUSE_DECL (innerc), t);
+ gsi_insert_after (&gsi, assign_stmt, GSI_CONTINUE_LINKING);
+ }
+ }
}
t = fold_convert (itype, s0);
@@ -11051,6 +11216,7 @@ expand_omp_for (struct omp_region *regio
extract_omp_for_data (as_a <gomp_for *> (last_stmt (region->entry)),
&fd, loops);
region->sched_kind = fd.sched_kind;
+ region->sched_modifiers = fd.sched_modifiers;
gcc_assert (EDGE_COUNT (region->entry->succs) == 2);
BRANCH_EDGE (region->entry)->flags &= ~EDGE_ABNORMAL;
@@ -11101,10 +11267,27 @@ expand_omp_for (struct omp_region *regio
&& fd.sched_kind == OMP_CLAUSE_SCHEDULE_STATIC)
fd.chunk_size = integer_zero_node;
gcc_assert (fd.sched_kind != OMP_CLAUSE_SCHEDULE_AUTO);
- fn_index = (fd.sched_kind == OMP_CLAUSE_SCHEDULE_RUNTIME)
- ? 3 : fd.sched_kind;
+ switch (fd.sched_kind)
+ {
+ case OMP_CLAUSE_SCHEDULE_RUNTIME:
+ fn_index = 3;
+ break;
+ case OMP_CLAUSE_SCHEDULE_DYNAMIC:
+ case OMP_CLAUSE_SCHEDULE_GUIDED:
+ if ((fd.sched_modifiers & OMP_CLAUSE_SCHEDULE_NONMONOTONIC)
+ && !fd.ordered
+ && !fd.have_ordered)
+ {
+ fn_index = 3 + fd.sched_kind;
+ break;
+ }
+ /* FALLTHRU */
+ default:
+ fn_index = fd.sched_kind;
+ break;
+ }
if (!fd.ordered)
- fn_index += fd.have_ordered * 4;
+ fn_index += fd.have_ordered * 6;
if (fd.ordered)
start_ix = ((int)BUILT_IN_GOMP_LOOP_DOACROSS_STATIC_START) + fn_index;
else
@@ -12287,9 +12470,14 @@ expand_omp_target (struct omp_region *re
vec_safe_push (offload_funcs, child_fn);
#endif
+ bool need_asm = DECL_ASSEMBLER_NAME_SET_P (current_function_decl)
+ && !DECL_ASSEMBLER_NAME_SET_P (child_fn);
+
/* Fix the callgraph edges for child_cfun. Those for cfun will be
fixed in a following pass. */
push_cfun (child_cfun);
+ if (need_asm)
+ assign_assembler_name_if_neeeded (child_fn);
cgraph_edge::rebuild_edges ();
#ifdef ENABLE_OFFLOADING
@@ -12315,6 +12503,13 @@ expand_omp_target (struct omp_region *re
if (flag_checking && !loops_state_satisfies_p (LOOPS_NEED_FIXUP))
verify_loop_structure ();
pop_cfun ();
+
+ if (dump_file && !gimple_in_ssa_p (cfun))
+ {
+ omp_any_child_fn_dumped = true;
+ dump_function_header (dump_file, child_fn, dump_flags);
+ dump_function_to_file (child_fn, dump_file, dump_flags);
+ }
}
/* Emit a library call to launch the offloading region, or do data
@@ -12489,6 +12684,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:
{
@@ -12602,6 +12822,7 @@ expand_omp_target (struct omp_region *re
static void
expand_omp (struct omp_region *region)
{
+ omp_any_child_fn_dumped = false;
while (region)
{
location_t saved_location;
@@ -12685,6 +12906,12 @@ expand_omp (struct omp_region *region)
input_location = saved_location;
region = region->next;
}
+ if (omp_any_child_fn_dumped)
+ {
+ if (dump_file)
+ dump_function_header (dump_file, current_function_decl, dump_flags);
+ omp_any_child_fn_dumped = false;
+ }
}
@@ -13761,26 +13988,36 @@ lower_omp_for_lastprivate (struct omp_fo
&& TREE_CODE (n2) != INTEGER_CST
&& gimple_omp_for_combined_into_p (fd->for_stmt))
{
- struct omp_context *task_ctx = NULL;
+ struct omp_context *taskreg_ctx = NULL;
if (gimple_code (ctx->outer->stmt) == GIMPLE_OMP_FOR)
{
gomp_for *gfor = as_a <gomp_for *> (ctx->outer->stmt);
- if (gimple_omp_for_kind (gfor) == GF_OMP_FOR_KIND_FOR)
+ if (gimple_omp_for_kind (gfor) == GF_OMP_FOR_KIND_FOR
+ || gimple_omp_for_kind (gfor) == GF_OMP_FOR_KIND_DISTRIBUTE)
{
- struct omp_for_data outer_fd;
- extract_omp_for_data (gfor, &outer_fd, NULL);
- n2 = fold_convert (TREE_TYPE (n2), outer_fd.loop.n2);
+ if (gimple_omp_for_combined_into_p (gfor))
+ {
+ gcc_assert (ctx->outer->outer
+ && is_parallel_ctx (ctx->outer->outer));
+ taskreg_ctx = ctx->outer->outer;
+ }
+ else
+ {
+ struct omp_for_data outer_fd;
+ extract_omp_for_data (gfor, &outer_fd, NULL);
+ n2 = fold_convert (TREE_TYPE (n2), outer_fd.loop.n2);
+ }
}
else if (gimple_omp_for_kind (gfor) == GF_OMP_FOR_KIND_TASKLOOP)
- task_ctx = ctx->outer->outer;
+ taskreg_ctx = ctx->outer->outer;
}
- else if (is_task_ctx (ctx->outer))
- task_ctx = ctx->outer;
- if (task_ctx)
+ else if (is_taskreg_ctx (ctx->outer))
+ taskreg_ctx = ctx->outer;
+ if (taskreg_ctx)
{
int i;
tree innerc
- = find_omp_clause (gimple_omp_task_clauses (task_ctx->stmt),
+ = find_omp_clause (gimple_omp_taskreg_clauses (taskreg_ctx->stmt),
OMP_CLAUSE__LOOPTEMP_);
gcc_assert (innerc);
for (i = 0; i < fd->collapse; i++)
@@ -13794,7 +14031,7 @@ lower_omp_for_lastprivate (struct omp_fo
if (innerc)
n2 = fold_convert (TREE_TYPE (n2),
lookup_decl (OMP_CLAUSE_DECL (innerc),
- task_ctx));
+ taskreg_ctx));
}
}
cond = build2 (cond_code, boolean_type_node, fd->loop.v, n2);
@@ -14592,7 +14829,9 @@ lower_omp_target (gimple_stmt_iterator *
case GOMP_MAP_ALWAYS_FROM:
case GOMP_MAP_ALWAYS_TOFROM:
case GOMP_MAP_FIRSTPRIVATE_POINTER:
+ case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
case GOMP_MAP_STRUCT:
+ case GOMP_MAP_ALWAYS_POINTER:
break;
case GOMP_MAP_FORCE_ALLOC:
case GOMP_MAP_FORCE_TO:
@@ -14631,7 +14870,8 @@ lower_omp_target (gimple_stmt_iterator *
}
if (offloaded
- && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
{
if (TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
{
@@ -14650,12 +14890,6 @@ lower_omp_target (gimple_stmt_iterator *
continue;
}
- if (offloaded && OMP_CLAUSE_MAP_PRIVATE (c))
- {
- map_cnt++;
- continue;
- }
-
if (!maybe_lookup_field (var, ctx))
continue;
@@ -14808,7 +15042,9 @@ lower_omp_target (gimple_stmt_iterator *
nc = c;
ovar = OMP_CLAUSE_DECL (c);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+ || (OMP_CLAUSE_MAP_KIND (c)
+ == GOMP_MAP_FIRSTPRIVATE_REFERENCE)))
break;
if (!DECL_P (ovar))
{
@@ -14840,14 +15076,7 @@ lower_omp_target (gimple_stmt_iterator *
gcc_assert (DECL_P (ovar2));
ovar = ovar2;
}
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_PRIVATE (c))
- {
- if (!maybe_lookup_field ((splay_tree_key) &DECL_UID (ovar),
- ctx))
- continue;
- }
- else if (!maybe_lookup_field (ovar, ctx))
+ if (!maybe_lookup_field (ovar, ctx))
continue;
}
@@ -14857,12 +15086,7 @@ lower_omp_target (gimple_stmt_iterator *
if (nc)
{
var = lookup_decl_in_outer_ctx (ovar, ctx);
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_PRIVATE (c))
- x = build_sender_ref ((splay_tree_key) &DECL_UID (ovar),
- ctx);
- else
- x = build_sender_ref (ovar, ctx);
+ x = build_sender_ref (ovar, ctx);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
@@ -15341,7 +15565,7 @@ lower_omp_target (gimple_stmt_iterator *
}
break;
}
- /* Handle GOMP_MAP_FIRSTPRIVATE_POINTER in second pass,
+ /* Handle GOMP_MAP_FIRSTPRIVATE_{POINTER,REFERENCE} in second pass,
so that firstprivate vars holding OMP_CLAUSE_SIZE if needed
are already handled. */
for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
@@ -15351,7 +15575,8 @@ lower_omp_target (gimple_stmt_iterator *
default:
break;
case OMP_CLAUSE_MAP:
- if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
{
location_t clause_loc = OMP_CLAUSE_LOCATION (c);
HOST_WIDE_INT offset = 0;
@@ -15405,6 +15630,8 @@ lower_omp_target (gimple_stmt_iterator *
}
else
is_ref = is_reference (var);
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+ is_ref = false;
bool ref_to_array = false;
if (is_ref)
{
@@ -15456,8 +15683,10 @@ lower_omp_target (gimple_stmt_iterator *
else if (OMP_CLAUSE_CHAIN (c)
&& OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c))
== OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
- == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ && (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
+ == GOMP_MAP_FIRSTPRIVATE_POINTER
+ || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
+ == GOMP_MAP_FIRSTPRIVATE_REFERENCE)))
prev = c;
break;
case OMP_CLAUSE_PRIVATE:
@@ -16553,8 +16782,29 @@ simd_clone_clauses_extract (struct cgrap
int argno = TREE_INT_CST_LOW (decl);
if (OMP_CLAUSE_LINEAR_VARIABLE_STRIDE (t))
{
- clone_info->args[argno].arg_type
- = SIMD_CLONE_ARG_TYPE_LINEAR_VARIABLE_STEP;
+ enum cgraph_simd_clone_arg_type arg_type;
+ if (TREE_CODE (args[argno]) == REFERENCE_TYPE)
+ switch (OMP_CLAUSE_LINEAR_KIND (t))
+ {
+ case OMP_CLAUSE_LINEAR_REF:
+ arg_type
+ = SIMD_CLONE_ARG_TYPE_LINEAR_REF_VARIABLE_STEP;
+ break;
+ case OMP_CLAUSE_LINEAR_UVAL:
+ arg_type
+ = SIMD_CLONE_ARG_TYPE_LINEAR_UVAL_VARIABLE_STEP;
+ break;
+ case OMP_CLAUSE_LINEAR_VAL:
+ case OMP_CLAUSE_LINEAR_DEFAULT:
+ arg_type
+ = SIMD_CLONE_ARG_TYPE_LINEAR_VAL_VARIABLE_STEP;
+ break;
+ default:
+ gcc_unreachable ();
+ }
+ else
+ arg_type = SIMD_CLONE_ARG_TYPE_LINEAR_VARIABLE_STEP;
+ clone_info->args[argno].arg_type = arg_type;
clone_info->args[argno].linear_step = tree_to_shwi (step);
gcc_assert (clone_info->args[argno].linear_step >= 0
&& clone_info->args[argno].linear_step < n);
@@ -16732,7 +16982,19 @@ simd_clone_mangle (struct cgraph_node *n
}
break;
case SIMD_CLONE_ARG_TYPE_LINEAR_VARIABLE_STEP:
- pp_character (&pp, 's');
+ pp_string (&pp, "ls");
+ pp_unsigned_wide_integer (&pp, arg.linear_step);
+ break;
+ case SIMD_CLONE_ARG_TYPE_LINEAR_REF_VARIABLE_STEP:
+ pp_string (&pp, "Rs");
+ pp_unsigned_wide_integer (&pp, arg.linear_step);
+ break;
+ case SIMD_CLONE_ARG_TYPE_LINEAR_VAL_VARIABLE_STEP:
+ pp_string (&pp, "Ls");
+ pp_unsigned_wide_integer (&pp, arg.linear_step);
+ break;
+ case SIMD_CLONE_ARG_TYPE_LINEAR_UVAL_VARIABLE_STEP:
+ pp_string (&pp, "Us");
pp_unsigned_wide_integer (&pp, arg.linear_step);
break;
default:
@@ -16924,6 +17186,7 @@ simd_clone_adjust_argument_types (struct
adj.op = IPA_PARM_OP_COPY;
break;
case SIMD_CLONE_ARG_TYPE_LINEAR_UVAL_CONSTANT_STEP:
+ case SIMD_CLONE_ARG_TYPE_LINEAR_UVAL_VARIABLE_STEP:
if (node->definition)
node->simdclone->args[i].simd_array
= create_tmp_simd_array (IDENTIFIER_POINTER (DECL_NAME (parm)),
@@ -16932,6 +17195,7 @@ simd_clone_adjust_argument_types (struct
adj.op = IPA_PARM_OP_COPY;
break;
case SIMD_CLONE_ARG_TYPE_LINEAR_VAL_CONSTANT_STEP:
+ case SIMD_CLONE_ARG_TYPE_LINEAR_VAL_VARIABLE_STEP:
case SIMD_CLONE_ARG_TYPE_VECTOR:
if (INTEGRAL_TYPE_P (parm_type) || POINTER_TYPE_P (parm_type))
veclen = node->simdclone->vecsize_int;
@@ -17319,6 +17583,74 @@ ipa_simd_modify_function_body (struct cg
}
}
+/* Helper function of simd_clone_adjust, return linear step addend
+ of Ith argument. */
+
+static tree
+simd_clone_linear_addend (struct cgraph_node *node, unsigned int i,
+ tree addtype, basic_block entry_bb)
+{
+ tree ptype = NULL_TREE;
+ switch (node->simdclone->args[i].arg_type)
+ {
+ case SIMD_CLONE_ARG_TYPE_LINEAR_CONSTANT_STEP:
+ case SIMD_CLONE_ARG_TYPE_LINEAR_REF_CONSTANT_STEP:
+ case SIMD_CLONE_ARG_TYPE_LINEAR_VAL_CONSTANT_STEP:
+ case SIMD_CLONE_ARG_TYPE_LINEAR_UVAL_CONSTANT_STEP:
+ return build_int_cst (addtype, node->simdclone->args[i].linear_step);
+ case SIMD_CLONE_ARG_TYPE_LINEAR_VARIABLE_STEP:
+ case SIMD_CLONE_ARG_TYPE_LINEAR_REF_VARIABLE_STEP:
+ ptype = TREE_TYPE (node->simdclone->args[i].orig_arg);
+ break;
+ case SIMD_CLONE_ARG_TYPE_LINEAR_VAL_VARIABLE_STEP:
+ case SIMD_CLONE_ARG_TYPE_LINEAR_UVAL_VARIABLE_STEP:
+ ptype = TREE_TYPE (TREE_TYPE (node->simdclone->args[i].orig_arg));
+ break;
+ default:
+ gcc_unreachable ();
+ }
+
+ unsigned int idx = node->simdclone->args[i].linear_step;
+ tree arg = node->simdclone->args[idx].orig_arg;
+ gcc_assert (is_gimple_reg_type (TREE_TYPE (arg)));
+ gimple_stmt_iterator gsi = gsi_after_labels (entry_bb);
+ gimple *g;
+ tree ret;
+ if (is_gimple_reg (arg))
+ ret = get_or_create_ssa_default_def (cfun, arg);
+ else
+ {
+ g = gimple_build_assign (make_ssa_name (TREE_TYPE (arg)), arg);
+ gsi_insert_before (&gsi, g, GSI_SAME_STMT);
+ ret = gimple_assign_lhs (g);
+ }
+ if (TREE_CODE (TREE_TYPE (arg)) == REFERENCE_TYPE)
+ {
+ g = gimple_build_assign (make_ssa_name (TREE_TYPE (TREE_TYPE (arg))),
+ build_simple_mem_ref (ret));
+ gsi_insert_before (&gsi, g, GSI_SAME_STMT);
+ ret = gimple_assign_lhs (g);
+ }
+ if (!useless_type_conversion_p (addtype, TREE_TYPE (ret)))
+ {
+ g = gimple_build_assign (make_ssa_name (addtype), NOP_EXPR, ret);
+ gsi_insert_before (&gsi, g, GSI_SAME_STMT);
+ ret = gimple_assign_lhs (g);
+ }
+ if (POINTER_TYPE_P (ptype))
+ {
+ tree size = TYPE_SIZE_UNIT (TREE_TYPE (ptype));
+ if (size && TREE_CODE (size) == INTEGER_CST)
+ {
+ g = gimple_build_assign (make_ssa_name (addtype), MULT_EXPR,
+ ret, fold_convert (addtype, size));
+ gsi_insert_before (&gsi, g, GSI_SAME_STMT);
+ ret = gimple_assign_lhs (g);
+ }
+ }
+ return ret;
+}
+
/* Adjust the argument types in NODE to their appropriate vector
counterparts. */
@@ -17558,7 +17890,11 @@ simd_clone_adjust (struct cgraph_node *n
else if ((node->simdclone->args[i].arg_type
== SIMD_CLONE_ARG_TYPE_LINEAR_CONSTANT_STEP)
|| (node->simdclone->args[i].arg_type
- == SIMD_CLONE_ARG_TYPE_LINEAR_REF_CONSTANT_STEP))
+ == SIMD_CLONE_ARG_TYPE_LINEAR_REF_CONSTANT_STEP)
+ || (node->simdclone->args[i].arg_type
+ == SIMD_CLONE_ARG_TYPE_LINEAR_VARIABLE_STEP)
+ || (node->simdclone->args[i].arg_type
+ == SIMD_CLONE_ARG_TYPE_LINEAR_REF_VARIABLE_STEP))
{
tree orig_arg = node->simdclone->args[i].orig_arg;
gcc_assert (INTEGRAL_TYPE_P (TREE_TYPE (orig_arg))
@@ -17593,10 +17929,10 @@ simd_clone_adjust (struct cgraph_node *n
? PLUS_EXPR : POINTER_PLUS_EXPR;
tree addtype = INTEGRAL_TYPE_P (TREE_TYPE (orig_arg))
? TREE_TYPE (orig_arg) : sizetype;
- tree addcst
- = build_int_cst (addtype, node->simdclone->args[i].linear_step);
- g = gimple_build_assign (iter2, code, iter1, addcst);
+ tree addcst = simd_clone_linear_addend (node, i, addtype,
+ entry_bb);
gsi = gsi_last_bb (incr_bb);
+ g = gimple_build_assign (iter2, code, iter1, addcst);
gsi_insert_before (&gsi, g, GSI_SAME_STMT);
imm_use_iterator iter;
@@ -17618,7 +17954,9 @@ simd_clone_adjust (struct cgraph_node *n
}
}
else if (node->simdclone->args[i].arg_type
- == SIMD_CLONE_ARG_TYPE_LINEAR_UVAL_CONSTANT_STEP)
+ == SIMD_CLONE_ARG_TYPE_LINEAR_UVAL_CONSTANT_STEP
+ || (node->simdclone->args[i].arg_type
+ == SIMD_CLONE_ARG_TYPE_LINEAR_UVAL_VARIABLE_STEP))
{
tree orig_arg = node->simdclone->args[i].orig_arg;
tree def = ssa_default_def (cfun, orig_arg);
@@ -17655,8 +17993,8 @@ simd_clone_adjust (struct cgraph_node *n
? PLUS_EXPR : POINTER_PLUS_EXPR;
tree addtype = INTEGRAL_TYPE_P (TREE_TYPE (iter3))
? TREE_TYPE (iter3) : sizetype;
- tree addcst
- = build_int_cst (addtype, node->simdclone->args[i].linear_step);
+ tree addcst = simd_clone_linear_addend (node, i, addtype,
+ entry_bb);
g = gimple_build_assign (iter5, code, iter4, addcst);
gsi = gsi_last_bb (incr_bb);
gsi_insert_before (&gsi, g, GSI_SAME_STMT);
@@ -1262,12 +1262,16 @@ extern tree c_finish_omp_critical (locat
extern tree c_finish_omp_ordered (location_t, tree, tree);
extern void c_finish_omp_barrier (location_t);
extern tree c_finish_omp_atomic (location_t, enum tree_code, enum tree_code,
- tree, tree, tree, tree, tree, bool, bool);
+ tree, tree, tree, tree, tree, bool, bool,
+ bool = false);
extern void c_finish_omp_flush (location_t);
extern void c_finish_omp_taskwait (location_t);
extern void c_finish_omp_taskyield (location_t);
extern tree c_finish_omp_for (location_t, enum tree_code, tree, tree, tree,
tree, tree, tree, tree);
+extern bool c_omp_check_loop_iv (tree, tree, walk_tree_lh);
+extern bool c_omp_check_loop_iv_exprs (location_t, tree, tree, tree, tree,
+ walk_tree_lh);
extern tree c_finish_oacc_wait (location_t, tree, tree);
extern tree c_oacc_split_loop_clauses (tree, tree *);
extern void c_omp_split_clauses (location_t, enum tree_code, omp_clause_mask,
@@ -175,12 +175,14 @@ c_finish_omp_taskyield (location_t loc)
LOC is the location of the atomic statement. The value returned
is either error_mark_node (if the construct was erroneous) or an
OMP_ATOMIC* node which should be added to the current statement
- tree with add_stmt. */
+ tree with add_stmt. If TEST is set, avoid calling save_expr
+ or create_tmp_var*. */
tree
c_finish_omp_atomic (location_t loc, enum tree_code code,
enum tree_code opcode, tree lhs, tree rhs,
- tree v, tree lhs1, tree rhs1, bool swapped, bool seq_cst)
+ tree v, tree lhs1, tree rhs1, bool swapped, bool seq_cst,
+ bool test)
{
tree x, type, addr, pre = NULL_TREE;
@@ -212,8 +214,10 @@ c_finish_omp_atomic (location_t loc, enu
addr = build_unary_op (loc, ADDR_EXPR, lhs, 0);
if (addr == error_mark_node)
return error_mark_node;
- addr = save_expr (addr);
- if (TREE_CODE (addr) != SAVE_EXPR
+ if (!test)
+ addr = save_expr (addr);
+ if (!test
+ && TREE_CODE (addr) != SAVE_EXPR
&& (TREE_CODE (addr) != ADDR_EXPR
|| !VAR_P (TREE_OPERAND (addr, 0))))
{
@@ -269,12 +273,15 @@ c_finish_omp_atomic (location_t loc, enu
if (rhs1
&& VAR_P (rhs1)
&& VAR_P (lhs)
- && rhs1 != lhs)
+ && rhs1 != lhs
+ && !test)
{
if (code == OMP_ATOMIC)
- error_at (loc, "%<#pragma omp atomic update%> uses two different variables for memory");
+ error_at (loc, "%<#pragma omp atomic update%> uses two different "
+ "variables for memory");
else
- error_at (loc, "%<#pragma omp atomic capture%> uses two different variables for memory");
+ error_at (loc, "%<#pragma omp atomic capture%> uses two different "
+ "variables for memory");
return error_mark_node;
}
@@ -284,9 +291,10 @@ c_finish_omp_atomic (location_t loc, enu
location, just diagnose different variables. */
if (lhs1 && VAR_P (lhs1) && VAR_P (lhs))
{
- if (lhs1 != lhs)
+ if (lhs1 != lhs && !test)
{
- error_at (loc, "%<#pragma omp atomic capture%> uses two different variables for memory");
+ error_at (loc, "%<#pragma omp atomic capture%> uses two "
+ "different variables for memory");
return error_mark_node;
}
}
@@ -308,7 +316,8 @@ c_finish_omp_atomic (location_t loc, enu
x = omit_one_operand_loc (loc, type, x, lhs1addr);
else
{
- x = save_expr (x);
+ if (!test)
+ x = save_expr (x);
x = omit_two_operands_loc (loc, type, x, x, lhs1addr);
}
}
@@ -683,19 +692,177 @@ c_finish_omp_for (location_t locus, enum
OMP_FOR_INCR (t) = incrv;
OMP_FOR_BODY (t) = body;
OMP_FOR_PRE_BODY (t) = pre_body;
- if (code == OMP_FOR)
- OMP_FOR_ORIG_DECLS (t) = orig_declv;
+ OMP_FOR_ORIG_DECLS (t) = orig_declv;
SET_EXPR_LOCATION (t, locus);
- return add_stmt (t);
+ return t;
}
}
+/* Type for passing data in between c_omp_check_loop_iv and
+ c_omp_check_loop_iv_r. */
+
+struct c_omp_check_loop_iv_data
+{
+ tree declv;
+ bool fail;
+ location_t stmt_loc;
+ location_t expr_loc;
+ int kind;
+ walk_tree_lh lh;
+ hash_set<tree> *ppset;
+};
+
+/* Helper function called via walk_tree, to diagnose uses
+ of associated loop IVs inside of lb, b and incr expressions
+ of OpenMP loops. */
+
+static tree
+c_omp_check_loop_iv_r (tree *tp, int *walk_subtrees, void *data)
+{
+ struct c_omp_check_loop_iv_data *d
+ = (struct c_omp_check_loop_iv_data *) data;
+ if (DECL_P (*tp))
+ {
+ int i;
+ for (i = 0; i < TREE_VEC_LENGTH (d->declv); i++)
+ if (*tp == TREE_VEC_ELT (d->declv, i))
+ {
+ location_t loc = d->expr_loc;
+ if (loc == UNKNOWN_LOCATION)
+ loc = d->stmt_loc;
+ switch (d->kind)
+ {
+ case 0:
+ error_at (loc, "initializer expression refers to "
+ "iteration variable %qD", *tp);
+ break;
+ case 1:
+ error_at (loc, "condition expression refers to "
+ "iteration variable %qD", *tp);
+ break;
+ case 2:
+ error_at (loc, "increment expression refers to "
+ "iteration variable %qD", *tp);
+ break;
+ }
+ d->fail = true;
+ }
+ }
+ /* Don't walk dtors added by C++ wrap_cleanups_r. */
+ else if (TREE_CODE (*tp) == TRY_CATCH_EXPR
+ && TRY_CATCH_IS_CLEANUP (*tp))
+ {
+ *walk_subtrees = 0;
+ return walk_tree_1 (&TREE_OPERAND (*tp, 0), c_omp_check_loop_iv_r, data,
+ d->ppset, d->lh);
+ }
+
+ return NULL_TREE;
+}
+
+/* Diagnose invalid references to loop iterators in lb, b and incr
+ expressions. */
+
+bool
+c_omp_check_loop_iv (tree stmt, tree declv, walk_tree_lh lh)
+{
+ hash_set<tree> pset;
+ struct c_omp_check_loop_iv_data data;
+ int i;
+
+ data.declv = declv;
+ data.fail = false;
+ data.stmt_loc = EXPR_LOCATION (stmt);
+ data.lh = lh;
+ data.ppset = &pset;
+ for (i = 0; i < TREE_VEC_LENGTH (OMP_FOR_INIT (stmt)); i++)
+ {
+ tree init = TREE_VEC_ELT (OMP_FOR_INIT (stmt), i);
+ gcc_assert (TREE_CODE (init) == MODIFY_EXPR);
+ tree decl = TREE_OPERAND (init, 0);
+ tree cond = TREE_VEC_ELT (OMP_FOR_COND (stmt), i);
+ gcc_assert (COMPARISON_CLASS_P (cond));
+ gcc_assert (TREE_OPERAND (cond, 0) == decl);
+ tree incr = TREE_VEC_ELT (OMP_FOR_INCR (stmt), i);
+ data.expr_loc = EXPR_LOCATION (TREE_OPERAND (init, 1));
+ data.kind = 0;
+ walk_tree_1 (&TREE_OPERAND (init, 1),
+ c_omp_check_loop_iv_r, &data, &pset, lh);
+ /* Don't warn for C++ random access iterators here, the
+ expression then involves the subtraction and always refers
+ to the original value. The C++ FE needs to warn on those
+ earlier. */
+ if (decl == TREE_VEC_ELT (declv, i))
+ {
+ data.expr_loc = EXPR_LOCATION (cond);
+ data.kind = 1;
+ walk_tree_1 (&TREE_OPERAND (cond, 1),
+ c_omp_check_loop_iv_r, &data, &pset, lh);
+ }
+ if (TREE_CODE (incr) == MODIFY_EXPR)
+ {
+ gcc_assert (TREE_OPERAND (incr, 0) == decl);
+ incr = TREE_OPERAND (incr, 1);
+ data.kind = 2;
+ if (TREE_CODE (incr) == PLUS_EXPR
+ && TREE_OPERAND (incr, 1) == decl)
+ {
+ data.expr_loc = EXPR_LOCATION (TREE_OPERAND (incr, 0));
+ walk_tree_1 (&TREE_OPERAND (incr, 0),
+ c_omp_check_loop_iv_r, &data, &pset, lh);
+ }
+ else
+ {
+ data.expr_loc = EXPR_LOCATION (TREE_OPERAND (incr, 1));
+ walk_tree_1 (&TREE_OPERAND (incr, 1),
+ c_omp_check_loop_iv_r, &data, &pset, lh);
+ }
+ }
+ }
+ return !data.fail;
+}
+
+/* Similar, but allows to check the init or cond expressions individually. */
+
+bool
+c_omp_check_loop_iv_exprs (location_t stmt_loc, tree declv, tree decl,
+ tree init, tree cond, walk_tree_lh lh)
+{
+ hash_set<tree> pset;
+ struct c_omp_check_loop_iv_data data;
+
+ data.declv = declv;
+ data.fail = false;
+ data.stmt_loc = stmt_loc;
+ data.lh = lh;
+ data.ppset = &pset;
+ if (init)
+ {
+ data.expr_loc = EXPR_LOCATION (init);
+ data.kind = 0;
+ walk_tree_1 (&init,
+ c_omp_check_loop_iv_r, &data, &pset, lh);
+ }
+ if (cond)
+ {
+ gcc_assert (COMPARISON_CLASS_P (cond));
+ data.expr_loc = EXPR_LOCATION (init);
+ data.kind = 1;
+ if (TREE_OPERAND (cond, 0) == decl)
+ walk_tree_1 (&TREE_OPERAND (cond, 1),
+ c_omp_check_loop_iv_r, &data, &pset, lh);
+ else
+ walk_tree_1 (&TREE_OPERAND (cond, 0),
+ c_omp_check_loop_iv_r, &data, &pset, lh);
+ }
+ return !data.fail;
+}
+
/* This function splits clauses for OpenACC combined loop
constructs. OpenACC combined loop constructs are:
#pragma acc kernels loop
- #pragma acc parallel loop
-*/
+ #pragma acc parallel loop */
tree
c_oacc_split_loop_clauses (tree clauses, tree *not_loop_clauses)
@@ -972,10 +1139,24 @@ c_omp_split_clauses (location_t loc, enu
s = C_OMP_CLAUSE_SPLIT_FOR;
}
break;
- /* Lastprivate is allowed on for, sections and simd. In
+ /* Lastprivate is allowed on distribute, for, sections and simd. In
parallel {for{, simd},sections} we actually want to put it on
parallel rather than for or sections. */
case OMP_CLAUSE_LASTPRIVATE:
+ if (code == OMP_DISTRIBUTE)
+ {
+ s = C_OMP_CLAUSE_SPLIT_DISTRIBUTE;
+ break;
+ }
+ if ((mask & (OMP_CLAUSE_MASK_1
+ << PRAGMA_OMP_CLAUSE_DIST_SCHEDULE)) != 0)
+ {
+ c = build_omp_clause (OMP_CLAUSE_LOCATION (clauses),
+ OMP_CLAUSE_LASTPRIVATE);
+ OMP_CLAUSE_DECL (c) = OMP_CLAUSE_DECL (clauses);
+ OMP_CLAUSE_CHAIN (c) = cclauses[C_OMP_CLAUSE_SPLIT_DISTRIBUTE];
+ cclauses[C_OMP_CLAUSE_SPLIT_DISTRIBUTE] = c;
+ }
if (code == OMP_FOR || code == OMP_SECTIONS)
{
if ((mask & (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NUM_THREADS))
@@ -1212,6 +1393,23 @@ c_omp_declare_simd_clauses_to_numbers (t
continue;
}
OMP_CLAUSE_DECL (c) = build_int_cst (integer_type_node, idx);
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LINEAR
+ && OMP_CLAUSE_LINEAR_VARIABLE_STRIDE (c))
+ {
+ decl = OMP_CLAUSE_LINEAR_STEP (c);
+ for (arg = parms, idx = 0; arg;
+ arg = TREE_CHAIN (arg), idx++)
+ if (arg == decl)
+ break;
+ if (arg == NULL_TREE)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%qD is not an function argument", decl);
+ continue;
+ }
+ OMP_CLAUSE_LINEAR_STEP (c)
+ = build_int_cst (integer_type_node, idx);
+ }
}
clvec.safe_push (c);
}
@@ -1249,6 +1447,17 @@ c_omp_declare_simd_clauses_to_decls (tre
break;
gcc_assert (arg);
OMP_CLAUSE_DECL (c) = arg;
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LINEAR
+ && OMP_CLAUSE_LINEAR_VARIABLE_STRIDE (c))
+ {
+ idx = tree_to_shwi (OMP_CLAUSE_LINEAR_STEP (c));
+ for (arg = DECL_ARGUMENTS (fndecl), i = 0; arg;
+ arg = TREE_CHAIN (arg), i++)
+ if (i == idx)
+ break;
+ gcc_assert (arg);
+ OMP_CLAUSE_LINEAR_STEP (c) = arg;
+ }
}
}
@@ -60,6 +60,8 @@ along with GCC; see the file COPYING3.
#include "builtins.h"
#include "gomp-constants.h"
#include "c-family/c-indentation.h"
+#include "gimple-expr.h"
+#include "context.h"
/* Initialization routine for this file. */
@@ -11574,35 +11576,58 @@ c_parser_omp_clause_reduction (c_parser
OpenMP 4.5:
schedule ( schedule-modifier : schedule-kind )
- schedule ( schedule-modifier : schedule-kind , expression )
+ schedule ( schedule-modifier [ , schedule-modifier ] : schedule-kind , expression )
schedule-modifier:
- simd */
+ simd
+ monotonic
+ nonmonotonic */
static tree
c_parser_omp_clause_schedule (c_parser *parser, tree list)
{
tree c, t;
location_t loc = c_parser_peek_token (parser)->location;
+ int modifiers = 0, nmodifiers = 0;
if (!c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>"))
return list;
c = build_omp_clause (loc, OMP_CLAUSE_SCHEDULE);
- if (c_parser_next_token_is (parser, CPP_NAME))
+ while (c_parser_next_token_is (parser, CPP_NAME))
{
tree kind = c_parser_peek_token (parser)->value;
const char *p = IDENTIFIER_POINTER (kind);
- if (strcmp ("simd", p) == 0
- && c_parser_peek_2nd_token (parser)->type == CPP_COLON)
+ if (strcmp ("simd", p) == 0)
+ OMP_CLAUSE_SCHEDULE_SIMD (c) = 1;
+ else if (strcmp ("monotonic", p) == 0)
+ modifiers |= OMP_CLAUSE_SCHEDULE_MONOTONIC;
+ else if (strcmp ("nonmonotonic", p) == 0)
+ modifiers |= OMP_CLAUSE_SCHEDULE_NONMONOTONIC;
+ else
+ break;
+ c_parser_consume_token (parser);
+ if (nmodifiers++ == 0
+ && c_parser_next_token_is (parser, CPP_COMMA))
+ c_parser_consume_token (parser);
+ else
{
- OMP_CLAUSE_SCHEDULE_SIMD (c) = 1;
- c_parser_consume_token (parser);
- c_parser_consume_token (parser);
+ c_parser_require (parser, CPP_COLON, "expected %<:%>");
+ break;
}
}
+ if ((modifiers & (OMP_CLAUSE_SCHEDULE_MONOTONIC
+ | OMP_CLAUSE_SCHEDULE_NONMONOTONIC))
+ == (OMP_CLAUSE_SCHEDULE_MONOTONIC
+ | OMP_CLAUSE_SCHEDULE_NONMONOTONIC))
+ {
+ error_at (loc, "both %<monotonic%> and %<nonmonotonic%> modifiers "
+ "specified");
+ modifiers = 0;
+ }
+
if (c_parser_next_token_is (parser, CPP_NAME))
{
tree kind = c_parser_peek_token (parser)->value;
@@ -11668,6 +11693,10 @@ c_parser_omp_clause_schedule (c_parser *
c_parser_skip_until_found (parser, CPP_CLOSE_PAREN,
"expected %<,%> or %<)%>");
+ OMP_CLAUSE_SCHEDULE_KIND (c)
+ = (enum omp_clause_schedule_kind)
+ (OMP_CLAUSE_SCHEDULE_KIND (c) | modifiers);
+
check_no_duplicate_clause (list, OMP_CLAUSE_SCHEDULE, "schedule");
OMP_CLAUSE_CHAIN (c) = list;
return c;
@@ -14139,8 +14168,15 @@ c_parser_omp_for_loop (location_t loc, c
{
stmt = c_finish_omp_for (loc, code, declv, NULL, initv, condv,
incrv, body, pre_body);
+
+ /* Check for iterators appearing in lb, b or incr expressions. */
+ if (stmt && !c_omp_check_loop_iv (stmt, declv, NULL))
+ stmt = NULL_TREE;
+
if (stmt)
{
+ add_stmt (stmt);
+
if (cclauses != NULL
&& cclauses[C_OMP_CLAUSE_SPLIT_PARALLEL] != NULL)
{
@@ -14823,6 +14859,7 @@ c_parser_omp_cancellation_point (c_parse
#define OMP_DISTRIBUTE_CLAUSE_MASK \
( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_FIRSTPRIVATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_LASTPRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DIST_SCHEDULE)\
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COLLAPSE))
@@ -14998,6 +15035,7 @@ c_parser_omp_target_data (location_t loc
map_seen = 3;
break;
case GOMP_MAP_FIRSTPRIVATE_POINTER:
+ case GOMP_MAP_ALWAYS_POINTER:
break;
default:
map_seen |= 1;
@@ -15131,6 +15169,7 @@ c_parser_omp_target_enter_data (location
map_seen = 3;
break;
case GOMP_MAP_FIRSTPRIVATE_POINTER:
+ case GOMP_MAP_ALWAYS_POINTER:
break;
default:
map_seen |= 1;
@@ -15217,6 +15256,7 @@ c_parser_omp_target_exit_data (location_
map_seen = 3;
break;
case GOMP_MAP_FIRSTPRIVATE_POINTER:
+ case GOMP_MAP_ALWAYS_POINTER:
break;
default:
map_seen |= 1;
@@ -15342,6 +15382,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];
@@ -15410,6 +15476,7 @@ check_clauses:
case GOMP_MAP_ALWAYS_TOFROM:
case GOMP_MAP_ALLOC:
case GOMP_MAP_FIRSTPRIVATE_POINTER:
+ case GOMP_MAP_ALWAYS_POINTER:
break;
default:
error_at (OMP_CLAUSE_LOCATION (*pc),
@@ -15679,6 +15746,7 @@ c_parser_omp_declare_target (c_parser *p
{
clauses = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_TO_DECLARE,
clauses);
+ clauses = c_finish_omp_clauses (clauses, true);
c_parser_skip_to_pragma_eol (parser);
}
else
@@ -15712,7 +15780,22 @@ c_parser_omp_declare_target (c_parser *p
continue;
}
if (!at1)
- DECL_ATTRIBUTES (t) = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (t));
+ {
+ symtab_node *node = symtab_node::get (t);
+ DECL_ATTRIBUTES (t) = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (t));
+ if (node != NULL)
+ {
+ node->offloadable = 1;
+#ifdef ENABLE_OFFLOADING
+ g->have_offload = true;
+ if (is_a <varpool_node *> (node))
+ {
+ vec_safe_push (offload_vars, t);
+ node->force_output = 1;
+ }
+#endif
+ }
+ }
}
}
@@ -4411,13 +4411,7 @@ c_decl_attributes (tree *node, tree attr
|| TREE_CODE (*node) == FUNCTION_DECL))
{
if (VAR_P (*node)
- && ((DECL_CONTEXT (*node)
- && TREE_CODE (DECL_CONTEXT (*node)) == FUNCTION_DECL)
- || (current_function_decl && !DECL_EXTERNAL (*node))))
- error ("%q+D in block scope inside of declare target directive",
- *node);
- else if (VAR_P (*node)
- && !lang_hooks.types.omp_mappable_type (TREE_TYPE (*node)))
+ && !lang_hooks.types.omp_mappable_type (TREE_TYPE (*node)))
error ("%q+D in declare target directive does not have mappable type",
*node);
else
@@ -11783,13 +11783,6 @@ handle_omp_array_sections_1 (tree c, tre
&& (TREE_CODE (length) != INTEGER_CST || integer_onep (length)))
first_non_one++;
}
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION
- && !integer_zerop (low_bound))
- {
- error_at (OMP_CLAUSE_LOCATION (c),
- "%<reduction%> array section has to be zero-based");
- return error_mark_node;
- }
if (TREE_CODE (type) == ARRAY_TYPE)
{
if (length == NULL_TREE
@@ -12130,7 +12123,24 @@ handle_omp_array_sections (tree c, bool
tree ptype = build_pointer_type (eltype);
if (TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE)
t = build_fold_addr_expr (t);
- t = build2 (MEM_REF, type, t, build_int_cst (ptype, 0));
+ tree t2 = build_fold_addr_expr (first);
+ t2 = fold_convert_loc (OMP_CLAUSE_LOCATION (c),
+ ptrdiff_type_node, t2);
+ t2 = fold_build2_loc (OMP_CLAUSE_LOCATION (c), MINUS_EXPR,
+ ptrdiff_type_node, t2,
+ fold_convert_loc (OMP_CLAUSE_LOCATION (c),
+ ptrdiff_type_node, t));
+ t2 = c_fully_fold (t2, false, NULL);
+ if (tree_fits_shwi_p (t2))
+ t = build2 (MEM_REF, type, t,
+ build_int_cst (ptype, tree_to_shwi (t2)));
+ else
+ {
+ t2 = fold_convert_loc (OMP_CLAUSE_LOCATION (c), sizetype, t2);
+ t = build2_loc (OMP_CLAUSE_LOCATION (c), POINTER_PLUS_EXPR,
+ TREE_TYPE (t), t, t2);
+ t = build2 (MEM_REF, type, t, build_int_cst (ptype, 0));
+ }
OMP_CLAUSE_DECL (c) = t;
return false;
}
@@ -12162,10 +12172,14 @@ handle_omp_array_sections (tree c, bool
break;
}
tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (c2, is_omp
- ? GOMP_MAP_FIRSTPRIVATE_POINTER
- : GOMP_MAP_POINTER);
- if (!is_omp && !c_mark_addressable (t))
+ if (!is_omp)
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER);
+ else if (TREE_CODE (t) == COMPONENT_REF)
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
+ else
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
+ if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
+ && !c_mark_addressable (t))
return false;
OMP_CLAUSE_DECL (c2) = t;
t = build_fold_addr_expr (first);
@@ -12233,12 +12247,15 @@ tree
c_finish_omp_clauses (tree clauses, bool is_omp, bool declare_simd)
{
bitmap_head generic_head, firstprivate_head, lastprivate_head;
- bitmap_head aligned_head, map_head, map_field_head, generic_field_head;
+ bitmap_head aligned_head, map_head, map_field_head;
tree c, t, type, *pc;
tree simdlen = NULL_TREE, safelen = NULL_TREE;
bool branch_seen = false;
bool copyprivate_seen = false;
+ bool linear_variable_step_check = false;
tree *nowait_clause = NULL;
+ bool ordered_seen = false;
+ tree schedule_clause = NULL_TREE;
bitmap_obstack_initialize (NULL);
bitmap_initialize (&generic_head, &bitmap_default_obstack);
@@ -12247,7 +12264,6 @@ c_finish_omp_clauses (tree clauses, bool
bitmap_initialize (&aligned_head, &bitmap_default_obstack);
bitmap_initialize (&map_head, &bitmap_default_obstack);
bitmap_initialize (&map_field_head, &bitmap_default_obstack);
- bitmap_initialize (&generic_field_head, &bitmap_default_obstack);
for (pc = &clauses, c = clauses; c ; c = *pc)
{
@@ -12468,6 +12484,8 @@ c_finish_omp_clauses (tree clauses, bool
break;
}
t = TREE_OPERAND (t, 0);
+ if (TREE_CODE (t) == POINTER_PLUS_EXPR)
+ t = TREE_OPERAND (t, 0);
if (TREE_CODE (t) == ADDR_EXPR)
t = TREE_OPERAND (t, 0);
}
@@ -12517,6 +12535,27 @@ c_finish_omp_clauses (tree clauses, bool
remove = true;
break;
}
+ if (declare_simd)
+ {
+ tree s = OMP_CLAUSE_LINEAR_STEP (c);
+ if (TREE_CODE (s) == PARM_DECL)
+ {
+ OMP_CLAUSE_LINEAR_VARIABLE_STRIDE (c) = 1;
+ /* map_head bitmap is used as uniform_head if
+ declare_simd. */
+ if (!bitmap_bit_p (&map_head, DECL_UID (s)))
+ linear_variable_step_check = true;
+ goto check_dup_generic;
+ }
+ if (TREE_CODE (s) != INTEGER_CST)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%<linear%> clause step %qE is neither constant "
+ "nor a parameter", s);
+ remove = true;
+ break;
+ }
+ }
if (TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL (c))) == POINTER_TYPE)
{
tree s = OMP_CLAUSE_LINEAR_STEP (c);
@@ -12553,6 +12592,12 @@ c_finish_omp_clauses (tree clauses, bool
"%qE appears more than once in data clauses", t);
remove = true;
}
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_PRIVATE
+ && bitmap_bit_p (&map_head, DECL_UID (t)))
+ {
+ error ("%qD appears both in data and map clauses", t);
+ remove = true;
+ }
else
bitmap_set_bit (&generic_head, DECL_UID (t));
break;
@@ -12574,6 +12619,11 @@ c_finish_omp_clauses (tree clauses, bool
"%qE appears more than once in data clauses", t);
remove = true;
}
+ else if (bitmap_bit_p (&map_head, DECL_UID (t)))
+ {
+ error ("%qD appears both in data and map clauses", t);
+ remove = true;
+ }
else
bitmap_set_bit (&firstprivate_head, DECL_UID (t));
break;
@@ -12767,14 +12817,7 @@ c_finish_omp_clauses (tree clauses, bool
break;
if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
{
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && (OMP_CLAUSE_MAP_KIND (c)
- == GOMP_MAP_FIRSTPRIVATE_POINTER))
- {
- if (bitmap_bit_p (&generic_field_head, DECL_UID (t)))
- break;
- }
- else if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
+ if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
break;
}
}
@@ -12817,13 +12860,13 @@ c_finish_omp_clauses (tree clauses, bool
error ("%qD appears more than once in data clauses", t);
remove = true;
}
- else
+ else if (bitmap_bit_p (&map_head, DECL_UID (t)))
{
- bitmap_set_bit (&generic_head, DECL_UID (t));
- if (t != OMP_CLAUSE_DECL (c)
- && TREE_CODE (OMP_CLAUSE_DECL (c)) == COMPONENT_REF)
- bitmap_set_bit (&generic_field_head, DECL_UID (t));
+ error ("%qD appears both in data and map clauses", t);
+ remove = true;
}
+ else
+ bitmap_set_bit (&generic_head, DECL_UID (t));
}
else if (bitmap_bit_p (&map_head, DECL_UID (t)))
{
@@ -12833,6 +12876,12 @@ c_finish_omp_clauses (tree clauses, bool
error ("%qD appears more than once in map clauses", t);
remove = true;
}
+ else if (bitmap_bit_p (&generic_head, DECL_UID (t))
+ || bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
+ {
+ error ("%qD appears both in data and map clauses", t);
+ remove = true;
+ }
else
{
bitmap_set_bit (&map_head, DECL_UID (t));
@@ -12843,17 +12892,22 @@ c_finish_omp_clauses (tree clauses, bool
break;
case OMP_CLAUSE_TO_DECLARE:
- t = OMP_CLAUSE_DECL (c);
- if (TREE_CODE (t) == FUNCTION_DECL)
- break;
- /* FALLTHRU */
case OMP_CLAUSE_LINK:
t = OMP_CLAUSE_DECL (c);
- if (!VAR_P (t))
+ if (TREE_CODE (t) == FUNCTION_DECL
+ && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO_DECLARE)
+ ;
+ else if (!VAR_P (t))
{
- error_at (OMP_CLAUSE_LOCATION (c),
- "%qE is not a variable in clause %qs", t,
- omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO_DECLARE)
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%qE is neither a variable nor a function name in "
+ "clause %qs", t,
+ omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+ else
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%qE is not a variable in clause %qs", t,
+ omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
else if (DECL_THREAD_LOCAL_P (t))
@@ -12870,6 +12924,17 @@ c_finish_omp_clauses (tree clauses, bool
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
+ if (remove)
+ break;
+ if (bitmap_bit_p (&generic_head, DECL_UID (t)))
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%qE appears more than once on the same "
+ "%<declare target%> directive", t);
+ remove = true;
+ }
+ else
+ bitmap_set_bit (&generic_head, DECL_UID (t));
break;
case OMP_CLAUSE_UNIFORM:
@@ -12885,6 +12950,8 @@ c_finish_omp_clauses (tree clauses, bool
remove = true;
break;
}
+ /* map_head bitmap is used as uniform_head if declare_simd. */
+ bitmap_set_bit (&map_head, DECL_UID (t));
goto check_dup_generic;
case OMP_CLAUSE_IS_DEVICE_PTR:
@@ -12917,8 +12984,6 @@ c_finish_omp_clauses (tree clauses, bool
case OMP_CLAUSE_NUM_THREADS:
case OMP_CLAUSE_NUM_TEAMS:
case OMP_CLAUSE_THREAD_LIMIT:
- case OMP_CLAUSE_SCHEDULE:
- case OMP_CLAUSE_ORDERED:
case OMP_CLAUSE_DEFAULT:
case OMP_CLAUSE_UNTIED:
case OMP_CLAUSE_COLLAPSE:
@@ -12953,6 +13018,39 @@ c_finish_omp_clauses (tree clauses, bool
pc = &OMP_CLAUSE_CHAIN (c);
continue;
+ case OMP_CLAUSE_SCHEDULE:
+ if (OMP_CLAUSE_SCHEDULE_KIND (c) & OMP_CLAUSE_SCHEDULE_NONMONOTONIC)
+ {
+ const char *p = NULL;
+ switch (OMP_CLAUSE_SCHEDULE_KIND (c) & OMP_CLAUSE_SCHEDULE_MASK)
+ {
+ case OMP_CLAUSE_SCHEDULE_STATIC: p = "static"; break;
+ case OMP_CLAUSE_SCHEDULE_DYNAMIC: break;
+ case OMP_CLAUSE_SCHEDULE_GUIDED: break;
+ case OMP_CLAUSE_SCHEDULE_AUTO: p = "auto"; break;
+ case OMP_CLAUSE_SCHEDULE_RUNTIME: p = "runtime"; break;
+ default: gcc_unreachable ();
+ }
+ if (p)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%<nonmonotonic%> modifier specified for %qs "
+ "schedule kind", p);
+ OMP_CLAUSE_SCHEDULE_KIND (c)
+ = (enum omp_clause_schedule_kind)
+ (OMP_CLAUSE_SCHEDULE_KIND (c)
+ & ~OMP_CLAUSE_SCHEDULE_NONMONOTONIC);
+ }
+ }
+ schedule_clause = c;
+ pc = &OMP_CLAUSE_CHAIN (c);
+ continue;
+
+ case OMP_CLAUSE_ORDERED:
+ ordered_seen = true;
+ pc = &OMP_CLAUSE_CHAIN (c);
+ continue;
+
case OMP_CLAUSE_SAFELEN:
safelen = c;
pc = &OMP_CLAUSE_CHAIN (c);
@@ -13043,6 +13141,42 @@ c_finish_omp_clauses (tree clauses, bool
= OMP_CLAUSE_SAFELEN_EXPR (safelen);
}
+ if (ordered_seen
+ && schedule_clause
+ && (OMP_CLAUSE_SCHEDULE_KIND (schedule_clause)
+ & OMP_CLAUSE_SCHEDULE_NONMONOTONIC))
+ {
+ error_at (OMP_CLAUSE_LOCATION (schedule_clause),
+ "%<nonmonotonic%> schedule modifier specified together "
+ "with %<ordered%> clause");
+ OMP_CLAUSE_SCHEDULE_KIND (schedule_clause)
+ = (enum omp_clause_schedule_kind)
+ (OMP_CLAUSE_SCHEDULE_KIND (schedule_clause)
+ & ~OMP_CLAUSE_SCHEDULE_NONMONOTONIC);
+ }
+
+ if (linear_variable_step_check)
+ for (pc = &clauses, c = clauses; c ; c = *pc)
+ {
+ bool remove = false;
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LINEAR
+ && OMP_CLAUSE_LINEAR_VARIABLE_STRIDE (c)
+ && !bitmap_bit_p (&map_head,
+ DECL_UID (OMP_CLAUSE_LINEAR_STEP (c))))
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%<linear%> clause step is a parameter %qD not "
+ "specified in %<uniform%> clause",
+ OMP_CLAUSE_LINEAR_STEP (c));
+ remove = true;
+ }
+
+ if (remove)
+ *pc = OMP_CLAUSE_CHAIN (c);
+ else
+ pc = &OMP_CLAUSE_CHAIN (c);
+ }
+
bitmap_obstack_release (NULL);
return clauses;
}
@@ -6324,7 +6324,7 @@ extern tree begin_omp_task (void);
extern tree finish_omp_task (tree, tree);
extern tree finish_omp_for (location_t, enum tree_code,
tree, tree, tree, tree, tree,
- tree, tree, tree);
+ tree, tree, vec<tree> *, tree);
extern void finish_omp_atomic (enum tree_code, enum tree_code,
tree, tree, tree, tree, tree,
bool);
@@ -6334,7 +6334,7 @@ extern void finish_omp_taskwait (void)
extern void finish_omp_taskyield (void);
extern void finish_omp_cancel (tree);
extern void finish_omp_cancellation_point (tree);
-extern tree omp_privatize_field (tree);
+extern tree omp_privatize_field (tree, bool);
extern tree begin_transaction_stmt (location_t, tree *, int);
extern void finish_transaction_stmt (tree, tree, int, tree);
extern tree build_transaction_expr (location_t, tree, int, tree);
@@ -43,6 +43,7 @@ along with GCC; see the file COPYING3.
#include "omp-low.h"
#include "gomp-constants.h"
#include "c-family/c-indentation.h"
+#include "context.h"
/* The lexer. */
@@ -30498,31 +30499,44 @@ cp_parser_omp_clause_reduction (cp_parse
OpenMP 4.5:
schedule ( schedule-modifier : schedule-kind )
- schedule ( schedule-modifier : schedule-kind , expression )
+ schedule ( schedule-modifier [ , schedule-modifier ] : schedule-kind , expression )
schedule-modifier:
- simd */
+ simd
+ monotonic
+ nonmonotonic */
static tree
cp_parser_omp_clause_schedule (cp_parser *parser, tree list, location_t location)
{
tree c, t;
+ int modifiers = 0, nmodifiers = 0;
if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
return list;
c = build_omp_clause (location, OMP_CLAUSE_SCHEDULE);
- if (cp_lexer_next_token_is (parser->lexer, CPP_NAME))
+ while (cp_lexer_next_token_is (parser->lexer, CPP_NAME))
{
tree id = cp_lexer_peek_token (parser->lexer)->u.value;
const char *p = IDENTIFIER_POINTER (id);
- if (strcmp ("simd", p) == 0
- && cp_lexer_nth_token_is (parser->lexer, 2, CPP_COLON))
+ if (strcmp ("simd", p) == 0)
+ OMP_CLAUSE_SCHEDULE_SIMD (c) = 1;
+ else if (strcmp ("monotonic", p) == 0)
+ modifiers |= OMP_CLAUSE_SCHEDULE_MONOTONIC;
+ else if (strcmp ("nonmonotonic", p) == 0)
+ modifiers |= OMP_CLAUSE_SCHEDULE_NONMONOTONIC;
+ else
+ break;
+ cp_lexer_consume_token (parser->lexer);
+ if (nmodifiers++ == 0
+ && cp_lexer_next_token_is (parser->lexer, CPP_COMMA))
+ cp_lexer_consume_token (parser->lexer);
+ else
{
- OMP_CLAUSE_SCHEDULE_SIMD (c) = 1;
- cp_lexer_consume_token (parser->lexer);
- cp_lexer_consume_token (parser->lexer);
+ cp_parser_require (parser, CPP_COLON, RT_COLON);
+ break;
}
}
@@ -30563,6 +30577,16 @@ cp_parser_omp_clause_schedule (cp_parser
goto invalid_kind;
cp_lexer_consume_token (parser->lexer);
+ if ((modifiers & (OMP_CLAUSE_SCHEDULE_MONOTONIC
+ | OMP_CLAUSE_SCHEDULE_NONMONOTONIC))
+ == (OMP_CLAUSE_SCHEDULE_MONOTONIC
+ | OMP_CLAUSE_SCHEDULE_NONMONOTONIC))
+ {
+ error_at (location, "both %<monotonic%> and %<nonmonotonic%> modifiers "
+ "specified");
+ modifiers = 0;
+ }
+
if (cp_lexer_next_token_is (parser->lexer, CPP_COMMA))
{
cp_token *token;
@@ -30588,6 +30612,10 @@ cp_parser_omp_clause_schedule (cp_parser
else if (!cp_parser_require (parser, CPP_CLOSE_PAREN, RT_COMMA_CLOSE_PAREN))
goto resync_fail;
+ OMP_CLAUSE_SCHEDULE_KIND (c)
+ = (enum omp_clause_schedule_kind)
+ (OMP_CLAUSE_SCHEDULE_KIND (c) | modifiers);
+
check_no_duplicate_clause (list, OMP_CLAUSE_SCHEDULE, "schedule", location);
OMP_CLAUSE_CHAIN (c) = list;
return c;
@@ -30780,7 +30808,7 @@ cp_parser_omp_clause_aligned (cp_parser
static tree
cp_parser_omp_clause_linear (cp_parser *parser, tree list,
- bool is_cilk_simd_fn)
+ bool is_cilk_simd_fn, bool declare_simd)
{
tree nlist, c, step = integer_one_node;
bool colon;
@@ -30824,7 +30852,30 @@ cp_parser_omp_clause_linear (cp_parser *
if (colon)
{
- step = cp_parser_expression (parser);
+ step = NULL_TREE;
+ if (declare_simd
+ && cp_lexer_next_token_is (parser->lexer, CPP_NAME)
+ && cp_lexer_nth_token_is (parser->lexer, 2, CPP_CLOSE_PAREN))
+ {
+ cp_token *token = cp_lexer_peek_token (parser->lexer);
+ cp_parser_parse_tentatively (parser);
+ step = cp_parser_id_expression (parser, /*template_p=*/false,
+ /*check_dependency_p=*/true,
+ /*template_p=*/NULL,
+ /*declarator_p=*/false,
+ /*optional_p=*/false);
+ if (step != error_mark_node)
+ step = cp_parser_lookup_name_simple (parser, step, token->location);
+ if (step == error_mark_node)
+ {
+ step = NULL_TREE;
+ cp_parser_abort_tentative_parse (parser);
+ }
+ else if (!cp_parser_parse_definitely (parser))
+ step = NULL_TREE;
+ }
+ if (!step)
+ step = cp_parser_expression (parser);
if (is_cilk_simd_fn && TREE_CODE (step) == PARM_DECL)
{
@@ -31511,7 +31562,6 @@ cp_parser_omp_all_clauses (cp_parser *pa
tree clauses = NULL;
bool first = true;
cp_token *token = NULL;
- bool cilk_simd_fn = false;
while (cp_lexer_next_token_is_not (parser->lexer, CPP_PRAGMA_EOL))
{
@@ -31725,9 +31775,15 @@ cp_parser_omp_all_clauses (cp_parser *pa
c_name = "aligned";
break;
case PRAGMA_OMP_CLAUSE_LINEAR:
- if (((mask >> PRAGMA_CILK_CLAUSE_VECTORLENGTH) & 1) != 0)
- cilk_simd_fn = true;
- clauses = cp_parser_omp_clause_linear (parser, clauses, cilk_simd_fn);
+ {
+ bool cilk_simd_fn = false, declare_simd = false;
+ if (((mask >> PRAGMA_CILK_CLAUSE_VECTORLENGTH) & 1) != 0)
+ cilk_simd_fn = true;
+ else if (((mask >> PRAGMA_OMP_CLAUSE_UNIFORM) & 1) != 0)
+ declare_simd = true;
+ clauses = cp_parser_omp_clause_linear (parser, clauses,
+ cilk_simd_fn, declare_simd);
+ }
c_name = "linear";
break;
case PRAGMA_OMP_CLAUSE_DEPEND:
@@ -32510,6 +32566,7 @@ cp_parser_omp_for_loop_init (cp_parser *
tree &this_pre_body,
vec<tree, va_gc> *for_block,
tree &init,
+ tree &orig_init,
tree &decl,
tree &real_decl)
{
@@ -32607,6 +32664,7 @@ cp_parser_omp_for_loop_init (cp_parser *
cp_finish_decl (decl, init, !is_non_constant_init,
asm_specification,
LOOKUP_ONLYCONVERTING);
+ orig_init = init;
if (CLASS_TYPE_P (TREE_TYPE (decl)))
{
vec_safe_push (for_block, this_pre_body);
@@ -32667,7 +32725,7 @@ cp_parser_omp_for_loop_init (cp_parser *
decl = cp_parser_lookup_name_simple (parser, name,
token->location);
if (TREE_CODE (decl) == FIELD_DECL)
- add_private_clause = omp_privatize_field (decl);
+ add_private_clause = omp_privatize_field (decl, false);
}
cp_parser_abort_tentative_parse (parser);
cp_parser_parse_tentatively (parser);
@@ -32684,6 +32742,7 @@ cp_parser_omp_for_loop_init (cp_parser *
cp_parser_parse_definitely (parser);
cp_parser_require (parser, CPP_EQ, RT_EQ);
rhs = cp_parser_assignment_expression (parser);
+ orig_init = rhs;
finish_expr_stmt (build_x_modify_expr (EXPR_LOCATION (rhs),
decl, NOP_EXPR,
rhs,
@@ -32713,13 +32772,14 @@ static tree
cp_parser_omp_for_loop (cp_parser *parser, enum tree_code code, tree clauses,
tree *cclauses)
{
- tree init, cond, incr, body, decl, pre_body = NULL_TREE, ret;
+ tree init, orig_init, cond, incr, body, decl, pre_body = NULL_TREE, ret;
tree real_decl, initv, condv, incrv, declv;
tree this_pre_body, cl, ordered_cl = NULL_TREE;
location_t loc_first;
bool collapse_err = false;
int i, collapse = 1, ordered = 0, count, nbraces = 0;
vec<tree, va_gc> *for_block = make_tree_vector ();
+ auto_vec<tree, 4> orig_inits;
for (cl = clauses; cl; cl = OMP_CLAUSE_CHAIN (cl))
if (OMP_CLAUSE_CODE (cl) == OMP_CLAUSE_COLLAPSE)
@@ -32786,13 +32846,13 @@ cp_parser_omp_for_loop (cp_parser *parse
if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
return NULL;
- init = decl = real_decl = NULL;
+ init = orig_init = decl = real_decl = NULL;
this_pre_body = push_stmt_list ();
add_private_clause
= cp_parser_omp_for_loop_init (parser, code,
this_pre_body, for_block,
- init, decl, real_decl);
+ init, orig_init, decl, real_decl);
cp_parser_require (parser, CPP_SEMICOLON, RT_SEMICOLON);
if (this_pre_body)
@@ -32924,6 +32984,11 @@ cp_parser_omp_for_loop (cp_parser *parse
TREE_VEC_ELT (initv, i) = init;
TREE_VEC_ELT (condv, i) = cond;
TREE_VEC_ELT (incrv, i) = incr;
+ if (orig_init)
+ {
+ orig_inits.safe_grow_cleared (i + 1);
+ orig_inits[i] = orig_init;
+ }
if (i == count - 1)
break;
@@ -32981,7 +33046,7 @@ cp_parser_omp_for_loop (cp_parser *parse
ret = NULL_TREE;
else
ret = finish_omp_for (loc_first, code, declv, NULL, initv, condv, incrv,
- body, pre_body, clauses);
+ body, pre_body, &orig_inits, clauses);
while (nbraces)
{
@@ -33626,6 +33691,7 @@ cp_parser_omp_cancellation_point (cp_par
#define OMP_DISTRIBUTE_CLAUSE_MASK \
( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_FIRSTPRIVATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_LASTPRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DIST_SCHEDULE)\
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COLLAPSE))
@@ -33821,6 +33887,8 @@ cp_parser_omp_target_data (cp_parser *pa
map_seen = 3;
break;
case GOMP_MAP_FIRSTPRIVATE_POINTER:
+ case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
+ case GOMP_MAP_ALWAYS_POINTER:
break;
default:
map_seen |= 1;
@@ -33912,6 +33980,8 @@ cp_parser_omp_target_enter_data (cp_pars
map_seen = 3;
break;
case GOMP_MAP_FIRSTPRIVATE_POINTER:
+ case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
+ case GOMP_MAP_ALWAYS_POINTER:
break;
default:
map_seen |= 1;
@@ -33999,6 +34069,8 @@ cp_parser_omp_target_exit_data (cp_parse
map_seen = 3;
break;
case GOMP_MAP_FIRSTPRIVATE_POINTER:
+ case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
+ case GOMP_MAP_ALWAYS_POINTER:
break;
default:
map_seen |= 1;
@@ -34169,6 +34241,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];
@@ -34235,6 +34334,8 @@ check_clauses:
case GOMP_MAP_ALWAYS_TOFROM:
case GOMP_MAP_ALLOC:
case GOMP_MAP_FIRSTPRIVATE_POINTER:
+ case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
+ case GOMP_MAP_ALWAYS_POINTER:
break;
default:
error_at (OMP_CLAUSE_LOCATION (*pc),
@@ -34766,11 +34867,12 @@ cp_parser_omp_declare_target (cp_parser
{
clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_TO_DECLARE,
clauses);
- cp_parser_skip_to_pragma_eol (parser, pragma_tok);
+ clauses = finish_omp_clauses (clauses, true);
+ cp_parser_require_pragma_eol (parser, pragma_tok);
}
else
{
- cp_parser_skip_to_pragma_eol (parser, pragma_tok);
+ cp_parser_require_pragma_eol (parser, pragma_tok);
scope_chain->omp_declare_target_attribute++;
return;
}
@@ -34800,7 +34902,22 @@ cp_parser_omp_declare_target (cp_parser
continue;
}
if (!at1)
- DECL_ATTRIBUTES (t) = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (t));
+ {
+ symtab_node *node = symtab_node::get (t);
+ DECL_ATTRIBUTES (t) = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (t));
+ if (node != NULL)
+ {
+ node->offloadable = 1;
+#ifdef ENABLE_OFFLOADING
+ g->have_offload = true;
+ if (is_a <varpool_node *> (node))
+ {
+ vec_safe_push (offload_vars, t);
+ node->force_output = 1;
+ }
+#endif
+ }
+ }
}
}
@@ -34837,7 +34954,7 @@ cp_parser_omp_end_declare_target (cp_par
cp_parser_skip_to_pragma_eol (parser, pragma_tok);
return;
}
- cp_parser_skip_to_pragma_eol (parser, pragma_tok);
+ cp_parser_require_pragma_eol (parser, pragma_tok);
if (!scope_chain->omp_declare_target_attribute)
error_at (pragma_tok->location,
"%<#pragma omp end declare target%> without corresponding "
@@ -1448,11 +1448,6 @@ cplus_decl_attributes (tree *decl, tree
&& DECL_CLASS_SCOPE_P (*decl))
error ("%q+D static data member inside of declare target directive",
*decl);
- else if (VAR_P (*decl)
- && (DECL_FUNCTION_SCOPE_P (*decl)
- || (current_function_decl && !DECL_EXTERNAL (*decl))))
- error ("%q+D in block scope inside of declare target directive",
- *decl);
else if (!processing_template_decl
&& VAR_P (*decl)
&& !cp_omp_mappable_type (TREE_TYPE (*decl)))
@@ -14419,7 +14419,6 @@ tsubst_omp_clauses (tree clauses, bool d
= tsubst_omp_clause_decl (OMP_CLAUSE_DECL (oc), args, complain,
in_decl);
break;
- case OMP_CLAUSE_LINEAR:
case OMP_CLAUSE_ALIGNED:
OMP_CLAUSE_DECL (nc)
= tsubst_omp_clause_decl (OMP_CLAUSE_DECL (oc), args, complain,
@@ -14427,12 +14426,25 @@ tsubst_omp_clauses (tree clauses, bool d
OMP_CLAUSE_OPERAND (nc, 1)
= tsubst_expr (OMP_CLAUSE_OPERAND (oc, 1), args, complain,
in_decl, /*integral_constant_expression_p=*/false);
- if (OMP_CLAUSE_CODE (oc) == OMP_CLAUSE_LINEAR
- && OMP_CLAUSE_LINEAR_STEP (oc) == NULL_TREE)
+ break;
+ case OMP_CLAUSE_LINEAR:
+ OMP_CLAUSE_DECL (nc)
+ = tsubst_omp_clause_decl (OMP_CLAUSE_DECL (oc), args, complain,
+ in_decl);
+ if (OMP_CLAUSE_LINEAR_STEP (oc) == NULL_TREE)
{
gcc_assert (!linear_no_step);
linear_no_step = nc;
}
+ else if (OMP_CLAUSE_LINEAR_VARIABLE_STRIDE (oc))
+ OMP_CLAUSE_LINEAR_STEP (nc)
+ = tsubst_omp_clause_decl (OMP_CLAUSE_LINEAR_STEP (oc), args,
+ complain, in_decl);
+ else
+ OMP_CLAUSE_LINEAR_STEP (nc)
+ = tsubst_expr (OMP_CLAUSE_LINEAR_STEP (oc), args, complain,
+ in_decl,
+ /*integral_constant_expression_p=*/false);
break;
case OMP_CLAUSE_NOWAIT:
case OMP_CLAUSE_DEFAULT:
@@ -14456,6 +14468,7 @@ tsubst_omp_clauses (tree clauses, bool d
if (allow_fields)
switch (OMP_CLAUSE_CODE (nc))
{
+ case OMP_CLAUSE_SHARED:
case OMP_CLAUSE_PRIVATE:
case OMP_CLAUSE_FIRSTPRIVATE:
case OMP_CLAUSE_LASTPRIVATE:
@@ -14635,7 +14648,7 @@ tsubst_omp_for_iterator (tree t, int i,
&& DECL_NAME (v) == this_identifier)
{
decl = TREE_OPERAND (decl, 1);
- decl = omp_privatize_field (decl);
+ decl = omp_privatize_field (decl, false);
}
/* FALLTHRU */
default:
@@ -14816,6 +14829,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. */
@@ -15240,7 +15274,7 @@ tsubst_expr (tree t, tree args, tsubst_f
if (OMP_FOR_INIT (t) != NULL_TREE)
{
declv = make_tree_vec (TREE_VEC_LENGTH (OMP_FOR_INIT (t)));
- if (TREE_CODE (t) == OMP_FOR && OMP_FOR_ORIG_DECLS (t))
+ if (OMP_FOR_ORIG_DECLS (t))
orig_declv = make_tree_vec (TREE_VEC_LENGTH (OMP_FOR_INIT (t)));
initv = make_tree_vec (TREE_VEC_LENGTH (OMP_FOR_INIT (t)));
condv = make_tree_vec (TREE_VEC_LENGTH (OMP_FOR_INIT (t)));
@@ -15267,7 +15301,7 @@ tsubst_expr (tree t, tree args, tsubst_f
if (OMP_FOR_INIT (t) != NULL_TREE)
t = finish_omp_for (EXPR_LOCATION (t), TREE_CODE (t), declv,
orig_declv, initv, condv, incrv, body, pre_body,
- clauses);
+ NULL, clauses);
else
{
t = make_node (TREE_CODE (t));
@@ -15318,6 +15352,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;
@@ -4323,19 +4323,20 @@ omp_note_field_privatization (tree f, tr
dummy VAR_DECL. */
tree
-omp_privatize_field (tree t)
+omp_privatize_field (tree t, bool shared)
{
tree m = finish_non_static_data_member (t, NULL_TREE, NULL_TREE);
if (m == error_mark_node)
return error_mark_node;
- if (!omp_private_member_map)
+ if (!omp_private_member_map && !shared)
omp_private_member_map = new hash_map<tree, tree>;
if (TREE_CODE (TREE_TYPE (t)) == REFERENCE_TYPE)
{
gcc_assert (TREE_CODE (m) == INDIRECT_REF);
m = TREE_OPERAND (m, 0);
}
- tree &v = omp_private_member_map->get_or_insert (t);
+ tree vb = NULL_TREE;
+ tree &v = shared ? vb : omp_private_member_map->get_or_insert (t);
if (v == NULL_TREE)
{
v = create_temporary_var (TREE_TYPE (m));
@@ -4344,7 +4345,8 @@ omp_privatize_field (tree t)
DECL_OMP_PRIVATIZED_MEMBER (v) = 1;
SET_DECL_VALUE_EXPR (v, m);
DECL_HAS_VALUE_EXPR_P (v) = 1;
- omp_private_member_vec.safe_push (t);
+ if (!shared)
+ omp_private_member_vec.safe_push (t);
}
return v;
}
@@ -4447,7 +4449,7 @@ handle_omp_array_sections_1 (tree c, tre
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION
&& TREE_CODE (TREE_CHAIN (t)) == FIELD_DECL)
- TREE_CHAIN (t) = omp_privatize_field (TREE_CHAIN (t));
+ TREE_CHAIN (t) = omp_privatize_field (TREE_CHAIN (t), false);
ret = handle_omp_array_sections_1 (c, TREE_CHAIN (t), types,
maybe_zero_len, first_non_one, is_omp);
if (ret == error_mark_node || ret == NULL_TREE)
@@ -4516,13 +4518,6 @@ handle_omp_array_sections_1 (tree c, tre
&& (TREE_CODE (length) != INTEGER_CST || integer_onep (length)))
first_non_one++;
}
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION
- && !integer_zerop (low_bound))
- {
- error_at (OMP_CLAUSE_LOCATION (c),
- "%<reduction%> array section has to be zero-based");
- return error_mark_node;
- }
if (TREE_CODE (type) == ARRAY_TYPE)
{
if (length == NULL_TREE
@@ -4863,7 +4858,24 @@ handle_omp_array_sections (tree c, bool
t = convert_from_reference (t);
else if (TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE)
t = build_fold_addr_expr (t);
- t = build2 (MEM_REF, type, t, build_int_cst (ptype, 0));
+ tree t2 = build_fold_addr_expr (first);
+ t2 = fold_convert_loc (OMP_CLAUSE_LOCATION (c),
+ ptrdiff_type_node, t2);
+ t2 = fold_build2_loc (OMP_CLAUSE_LOCATION (c), MINUS_EXPR,
+ ptrdiff_type_node, t2,
+ fold_convert_loc (OMP_CLAUSE_LOCATION (c),
+ ptrdiff_type_node, t));
+ if (tree_fits_shwi_p (t2))
+ t = build2 (MEM_REF, type, t,
+ build_int_cst (ptype, tree_to_shwi (t2)));
+ else
+ {
+ t2 = fold_convert_loc (OMP_CLAUSE_LOCATION (c),
+ sizetype, t2);
+ t = build2_loc (OMP_CLAUSE_LOCATION (c), POINTER_PLUS_EXPR,
+ TREE_TYPE (t), t, t2);
+ t = build2 (MEM_REF, type, t, build_int_cst (ptype, 0));
+ }
OMP_CLAUSE_DECL (c) = t;
return false;
}
@@ -4892,9 +4904,20 @@ handle_omp_array_sections (tree c, bool
}
tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (c2, is_omp ? GOMP_MAP_FIRSTPRIVATE_POINTER
- : GOMP_MAP_POINTER);
- if (!is_omp && !cxx_mark_addressable (t))
+ if (!is_omp)
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER);
+ else if (TREE_CODE (t) == COMPONENT_REF)
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
+ else if (REFERENCE_REF_P (t)
+ && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
+ {
+ t = TREE_OPERAND (t, 0);
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
+ }
+ else
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
+ if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
+ && !cxx_mark_addressable (t))
return false;
OMP_CLAUSE_DECL (c2) = t;
t = build_fold_addr_expr (first);
@@ -4912,15 +4935,18 @@ handle_omp_array_sections (tree c, bool
OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
OMP_CLAUSE_CHAIN (c) = c2;
ptr = OMP_CLAUSE_DECL (c2);
- if (!is_omp
+ if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
&& TREE_CODE (TREE_TYPE (ptr)) == REFERENCE_TYPE
&& POINTER_TYPE_P (TREE_TYPE (TREE_TYPE (ptr))))
{
tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (c3, GOMP_MAP_POINTER);
+ OMP_CLAUSE_SET_MAP_KIND (c3, OMP_CLAUSE_MAP_KIND (c2));
OMP_CLAUSE_DECL (c3) = ptr;
- OMP_CLAUSE_DECL (c2) = convert_from_reference (ptr);
+ if (OMP_CLAUSE_MAP_KIND (c2) == GOMP_MAP_ALWAYS_POINTER)
+ OMP_CLAUSE_DECL (c2) = build_simple_mem_ref (ptr);
+ else
+ OMP_CLAUSE_DECL (c2) = convert_from_reference (ptr);
OMP_CLAUSE_SIZE (c3) = size_zero_node;
OMP_CLAUSE_CHAIN (c3) = OMP_CLAUSE_CHAIN (c2);
OMP_CLAUSE_CHAIN (c2) = c3;
@@ -5645,11 +5671,12 @@ tree
finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd)
{
bitmap_head generic_head, firstprivate_head, lastprivate_head;
- bitmap_head aligned_head, map_head, map_field_head, generic_field_head;
+ bitmap_head aligned_head, map_head, map_field_head;
tree c, t, *pc;
tree safelen = NULL_TREE;
bool branch_seen = false;
bool copyprivate_seen = false;
+ bool ordered_seen = false;
bitmap_obstack_initialize (NULL);
bitmap_initialize (&generic_head, &bitmap_default_obstack);
@@ -5658,7 +5685,6 @@ finish_omp_clauses (tree clauses, bool a
bitmap_initialize (&aligned_head, &bitmap_default_obstack);
bitmap_initialize (&map_head, &bitmap_default_obstack);
bitmap_initialize (&map_field_head, &bitmap_default_obstack);
- bitmap_initialize (&generic_field_head, &bitmap_default_obstack);
for (pc = &clauses, c = clauses; c ; c = *pc)
{
@@ -5668,6 +5694,7 @@ finish_omp_clauses (tree clauses, bool a
switch (OMP_CLAUSE_CODE (c))
{
case OMP_CLAUSE_SHARED:
+ field_ok = allow_fields;
goto check_dup_generic;
case OMP_CLAUSE_PRIVATE:
field_ok = allow_fields;
@@ -5691,6 +5718,8 @@ finish_omp_clauses (tree clauses, bool a
{
gcc_assert (TREE_CODE (t) == MEM_REF);
t = TREE_OPERAND (t, 0);
+ if (TREE_CODE (t) == POINTER_PLUS_EXPR)
+ t = TREE_OPERAND (t, 0);
if (TREE_CODE (t) == ADDR_EXPR
|| TREE_CODE (t) == INDIRECT_REF)
t = TREE_OPERAND (t, 0);
@@ -5753,7 +5782,11 @@ finish_omp_clauses (tree clauses, bool a
break;
}
else if (!type_dependent_expression_p (t)
- && !INTEGRAL_TYPE_P (TREE_TYPE (t)))
+ && !INTEGRAL_TYPE_P (TREE_TYPE (t))
+ && (!declare_simd
+ || TREE_CODE (t) != PARM_DECL
+ || TREE_CODE (TREE_TYPE (t)) != REFERENCE_TYPE
+ || !INTEGRAL_TYPE_P (TREE_TYPE (TREE_TYPE (t)))))
{
error ("linear step expression must be integral");
remove = true;
@@ -5762,12 +5795,27 @@ finish_omp_clauses (tree clauses, bool a
else
{
t = mark_rvalue_use (t);
+ if (declare_simd && TREE_CODE (t) == PARM_DECL)
+ {
+ OMP_CLAUSE_LINEAR_VARIABLE_STRIDE (c) = 1;
+ goto check_dup_generic;
+ }
if (!processing_template_decl
&& (VAR_P (OMP_CLAUSE_DECL (c))
|| TREE_CODE (OMP_CLAUSE_DECL (c)) == PARM_DECL))
{
- if (TREE_CODE (OMP_CLAUSE_DECL (c)) == PARM_DECL)
- t = maybe_constant_value (t);
+ if (declare_simd)
+ {
+ t = maybe_constant_value (t);
+ if (TREE_CODE (t) != INTEGER_CST)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%<linear%> clause step %qE is neither "
+ "constant nor a parameter", t);
+ remove = true;
+ break;
+ }
+ }
t = fold_build_cleanup_point_expr (TREE_TYPE (t), t);
tree type = TREE_TYPE (OMP_CLAUSE_DECL (c));
if (TREE_CODE (type) == REFERENCE_TYPE)
@@ -5821,7 +5869,7 @@ finish_omp_clauses (tree clauses, bool a
t = omp_clause_decl_field (OMP_CLAUSE_DECL (c));
if (t)
{
- if (!remove)
+ if (!remove && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_SHARED)
omp_note_field_privatization (t, OMP_CLAUSE_DECL (c));
}
else
@@ -5857,6 +5905,12 @@ finish_omp_clauses (tree clauses, bool a
error ("%qD appears more than once in data clauses", t);
remove = true;
}
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_PRIVATE
+ && bitmap_bit_p (&map_head, DECL_UID (t)))
+ {
+ error ("%qD appears both in data and map clauses", t);
+ remove = true;
+ }
else
bitmap_set_bit (&generic_head, DECL_UID (t));
if (!field_ok)
@@ -5866,7 +5920,9 @@ finish_omp_clauses (tree clauses, bool a
&& TREE_CODE (t) == FIELD_DECL
&& t == OMP_CLAUSE_DECL (c))
{
- OMP_CLAUSE_DECL (c) = omp_privatize_field (t);
+ OMP_CLAUSE_DECL (c)
+ = omp_privatize_field (t, (OMP_CLAUSE_CODE (c)
+ == OMP_CLAUSE_SHARED));
if (OMP_CLAUSE_DECL (c) == error_mark_node)
remove = true;
}
@@ -5902,6 +5958,11 @@ finish_omp_clauses (tree clauses, bool a
error ("%qD appears more than once in data clauses", t);
remove = true;
}
+ else if (bitmap_bit_p (&map_head, DECL_UID (t)))
+ {
+ error ("%qD appears both in data and map clauses", t);
+ remove = true;
+ }
else
bitmap_set_bit (&firstprivate_head, DECL_UID (t));
goto handle_field_decl;
@@ -6072,6 +6133,30 @@ finish_omp_clauses (tree clauses, bool a
break;
case OMP_CLAUSE_SCHEDULE:
+ if (OMP_CLAUSE_SCHEDULE_KIND (c) & OMP_CLAUSE_SCHEDULE_NONMONOTONIC)
+ {
+ const char *p = NULL;
+ switch (OMP_CLAUSE_SCHEDULE_KIND (c) & OMP_CLAUSE_SCHEDULE_MASK)
+ {
+ case OMP_CLAUSE_SCHEDULE_STATIC: p = "static"; break;
+ case OMP_CLAUSE_SCHEDULE_DYNAMIC: break;
+ case OMP_CLAUSE_SCHEDULE_GUIDED: break;
+ case OMP_CLAUSE_SCHEDULE_AUTO: p = "auto"; break;
+ case OMP_CLAUSE_SCHEDULE_RUNTIME: p = "runtime"; break;
+ default: gcc_unreachable ();
+ }
+ if (p)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%<nonmonotonic%> modifier specified for %qs "
+ "schedule kind", p);
+ OMP_CLAUSE_SCHEDULE_KIND (c)
+ = (enum omp_clause_schedule_kind)
+ (OMP_CLAUSE_SCHEDULE_KIND (c)
+ & ~OMP_CLAUSE_SCHEDULE_NONMONOTONIC);
+ }
+ }
+
t = OMP_CLAUSE_SCHEDULE_CHUNK_EXPR (c);
if (t == NULL)
;
@@ -6406,7 +6491,10 @@ finish_omp_clauses (tree clauses, bool a
}
if (REFERENCE_REF_P (t)
&& TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
- t = TREE_OPERAND (t, 0);
+ {
+ t = TREE_OPERAND (t, 0);
+ OMP_CLAUSE_DECL (c) = t;
+ }
if (TREE_CODE (t) == COMPONENT_REF
&& allow_fields
&& OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_)
@@ -6443,15 +6531,8 @@ finish_omp_clauses (tree clauses, bool a
break;
if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
{
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && (OMP_CLAUSE_MAP_KIND (c)
- == GOMP_MAP_FIRSTPRIVATE_POINTER))
- {
- if (bitmap_bit_p (&generic_field_head, DECL_UID (t)))
- break;
- }
- else if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
- break;
+ if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
+ goto handle_map_references;
}
}
if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
@@ -6459,7 +6540,8 @@ finish_omp_clauses (tree clauses, bool a
if (processing_template_decl)
break;
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER)
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER))
break;
if (DECL_P (t))
error ("%qD is not a variable in %qs clause", t,
@@ -6511,17 +6593,13 @@ finish_omp_clauses (tree clauses, bool a
error ("%qD appears more than once in data clauses", t);
remove = true;
}
- else
+ else if (bitmap_bit_p (&map_head, DECL_UID (t)))
{
- bitmap_set_bit (&generic_head, DECL_UID (t));
- if (t != OMP_CLAUSE_DECL (c)
- && (TREE_CODE (OMP_CLAUSE_DECL (c)) == COMPONENT_REF
- || (REFERENCE_REF_P (OMP_CLAUSE_DECL (c))
- && (TREE_CODE (TREE_OPERAND (OMP_CLAUSE_DECL (c),
- 0))
- == COMPONENT_REF))))
- bitmap_set_bit (&generic_field_head, DECL_UID (t));
+ error ("%qD appears both in data and map clauses", t);
+ remove = true;
}
+ else
+ bitmap_set_bit (&generic_head, DECL_UID (t));
}
else if (bitmap_bit_p (&map_head, DECL_UID (t)))
{
@@ -6531,6 +6609,12 @@ finish_omp_clauses (tree clauses, bool a
error ("%qD appears more than once in map clauses", t);
remove = true;
}
+ else if (bitmap_bit_p (&generic_head, DECL_UID (t))
+ || bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
+ {
+ error ("%qD appears both in data and map clauses", t);
+ remove = true;
+ }
else
{
bitmap_set_bit (&map_head, DECL_UID (t));
@@ -6538,20 +6622,75 @@ finish_omp_clauses (tree clauses, bool a
&& TREE_CODE (OMP_CLAUSE_DECL (c)) == COMPONENT_REF)
bitmap_set_bit (&map_field_head, DECL_UID (t));
}
+ handle_map_references:
+ if (!remove
+ && !processing_template_decl
+ && allow_fields
+ && TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL (c))) == REFERENCE_TYPE)
+ {
+ t = OMP_CLAUSE_DECL (c);
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+ {
+ OMP_CLAUSE_DECL (c) = build_simple_mem_ref (t);
+ if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
+ OMP_CLAUSE_SIZE (c)
+ = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (t)));
+ }
+ else if (OMP_CLAUSE_MAP_KIND (c)
+ != GOMP_MAP_FIRSTPRIVATE_POINTER
+ && (OMP_CLAUSE_MAP_KIND (c)
+ != GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+ && (OMP_CLAUSE_MAP_KIND (c)
+ != GOMP_MAP_ALWAYS_POINTER))
+ {
+ tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+ OMP_CLAUSE_MAP);
+ if (TREE_CODE (t) == COMPONENT_REF)
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
+ else
+ OMP_CLAUSE_SET_MAP_KIND (c2,
+ GOMP_MAP_FIRSTPRIVATE_REFERENCE);
+ OMP_CLAUSE_DECL (c2) = t;
+ OMP_CLAUSE_SIZE (c2) = size_zero_node;
+ OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
+ OMP_CLAUSE_CHAIN (c) = c2;
+ OMP_CLAUSE_DECL (c) = build_simple_mem_ref (t);
+ if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
+ OMP_CLAUSE_SIZE (c)
+ = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (t)));
+ c = c2;
+ }
+ }
break;
case OMP_CLAUSE_TO_DECLARE:
- t = OMP_CLAUSE_DECL (c);
- if (TREE_CODE (t) == FUNCTION_DECL)
- break;
- /* FALLTHRU */
case OMP_CLAUSE_LINK:
t = OMP_CLAUSE_DECL (c);
- if (!VAR_P (t))
+ if (TREE_CODE (t) == FUNCTION_DECL
+ && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO_DECLARE)
+ ;
+ else if (!VAR_P (t))
{
- error_at (OMP_CLAUSE_LOCATION (c),
- "%qE is not a variable in clause %qs", t,
- omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO_DECLARE)
+ {
+ if (TREE_CODE (t) == OVERLOAD && OVL_CHAIN (t))
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "overloaded function name %qE in clause %qs", t,
+ omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+ else if (TREE_CODE (t) == TEMPLATE_ID_EXPR)
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "template %qE in clause %qs", t,
+ omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+ else
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%qE is neither a variable nor a function name "
+ "in clause %qs", t,
+ omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+ }
+ else
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%qE is not a variable in clause %qs", t,
+ omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
else if (DECL_THREAD_LOCAL_P (t))
@@ -6568,6 +6707,17 @@ finish_omp_clauses (tree clauses, bool a
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
+ if (remove)
+ break;
+ if (bitmap_bit_p (&generic_head, DECL_UID (t)))
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%qE appears more than once on the same "
+ "%<declare target%> directive", t);
+ remove = true;
+ }
+ else
+ bitmap_set_bit (&generic_head, DECL_UID (t));
break;
case OMP_CLAUSE_UNIFORM:
@@ -6583,6 +6733,8 @@ finish_omp_clauses (tree clauses, bool a
remove = true;
break;
}
+ /* map_head bitmap is used as uniform_head if declare_simd. */
+ bitmap_set_bit (&map_head, DECL_UID (t));
goto check_dup_generic;
case OMP_CLAUSE_GRAINSIZE:
@@ -6688,7 +6840,6 @@ finish_omp_clauses (tree clauses, bool a
goto check_dup_generic;
case OMP_CLAUSE_NOWAIT:
- case OMP_CLAUSE_ORDERED:
case OMP_CLAUSE_DEFAULT:
case OMP_CLAUSE_UNTIED:
case OMP_CLAUSE_COLLAPSE:
@@ -6707,6 +6858,10 @@ finish_omp_clauses (tree clauses, bool a
case OMP_CLAUSE_SEQ:
break;
+ case OMP_CLAUSE_ORDERED:
+ ordered_seen = true;
+ break;
+
case OMP_CLAUSE_INBRANCH:
case OMP_CLAUSE_NOTINBRANCH:
if (branch_seen)
@@ -6768,6 +6923,17 @@ finish_omp_clauses (tree clauses, bool a
case OMP_CLAUSE_LINEAR:
if (!declare_simd)
need_implicitly_determined = true;
+ else if (OMP_CLAUSE_LINEAR_VARIABLE_STRIDE (c)
+ && !bitmap_bit_p (&map_head,
+ DECL_UID (OMP_CLAUSE_LINEAR_STEP (c))))
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%<linear%> clause step is a parameter %qD not "
+ "specified in %<uniform%> clause",
+ OMP_CLAUSE_LINEAR_STEP (c));
+ *pc = OMP_CLAUSE_CHAIN (c);
+ continue;
+ }
break;
case OMP_CLAUSE_COPYPRIVATE:
need_copy_assignment = true;
@@ -6789,6 +6955,21 @@ finish_omp_clauses (tree clauses, bool a
}
pc = &OMP_CLAUSE_CHAIN (c);
continue;
+ case OMP_CLAUSE_SCHEDULE:
+ if (ordered_seen
+ && (OMP_CLAUSE_SCHEDULE_KIND (c)
+ & OMP_CLAUSE_SCHEDULE_NONMONOTONIC))
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%<nonmonotonic%> schedule modifier specified "
+ "together with %<ordered%> clause");
+ OMP_CLAUSE_SCHEDULE_KIND (c)
+ = (enum omp_clause_schedule_kind)
+ (OMP_CLAUSE_SCHEDULE_KIND (c)
+ & ~OMP_CLAUSE_SCHEDULE_NONMONOTONIC);
+ }
+ pc = &OMP_CLAUSE_CHAIN (c);
+ continue;
case OMP_CLAUSE_NOWAIT:
if (copyprivate_seen)
{
@@ -6906,6 +7087,15 @@ finish_omp_clauses (tree clauses, bool a
need_dtor))
remove = true;
+ if (!remove
+ && c_kind == OMP_CLAUSE_SHARED
+ && processing_template_decl)
+ {
+ t = omp_clause_decl_field (OMP_CLAUSE_DECL (c));
+ if (t)
+ OMP_CLAUSE_DECL (c) = t;
+ }
+
if (remove)
*pc = OMP_CLAUSE_CHAIN (c);
else
@@ -7202,9 +7392,10 @@ finish_omp_task (tree clauses, tree body
static bool
handle_omp_for_class_iterator (int i, location_t locus, enum tree_code code,
- tree declv, tree initv, tree condv, tree incrv,
- tree *body, tree *pre_body, tree &clauses,
- tree *lastp, int collapse, int ordered)
+ tree declv, tree orig_declv, tree initv,
+ tree condv, tree incrv, tree *body,
+ tree *pre_body, tree &clauses, tree *lastp,
+ int collapse, int ordered)
{
tree diff, iter_init, iter_incr = NULL, last;
tree incr_var = NULL, orig_pre_body, orig_body, c;
@@ -7261,6 +7452,10 @@ handle_omp_for_class_iterator (int i, lo
TREE_OPERAND (cond, 1), iter);
return true;
}
+ if (!c_omp_check_loop_iv_exprs (locus, orig_declv,
+ TREE_VEC_ELT (declv, i), NULL_TREE,
+ cond, cp_walk_subtrees))
+ return true;
switch (TREE_CODE (incr))
{
@@ -7516,7 +7711,7 @@ handle_omp_for_class_iterator (int i, lo
tree
finish_omp_for (location_t locus, enum tree_code code, tree declv,
tree orig_declv, tree initv, tree condv, tree incrv,
- tree body, tree pre_body, tree clauses)
+ tree body, tree pre_body, vec<tree> *orig_inits, tree clauses)
{
tree omp_for = NULL, orig_incr = NULL;
tree decl = NULL, init, cond, incr, orig_decl = NULL_TREE, block = NULL_TREE;
@@ -7592,6 +7787,20 @@ finish_omp_for (location_t locus, enum t
TREE_VEC_ELT (initv, i) = init;
}
+ if (orig_inits)
+ {
+ bool fail = false;
+ tree orig_init;
+ FOR_EACH_VEC_ELT (*orig_inits, i, orig_init)
+ if (orig_init
+ && !c_omp_check_loop_iv_exprs (locus, declv,
+ TREE_VEC_ELT (declv, i), orig_init,
+ NULL_TREE, cp_walk_subtrees))
+ fail = true;
+ if (fail)
+ return NULL;
+ }
+
if (dependent_omp_for_p (declv, initv, condv, incrv))
{
tree stmt;
@@ -7665,10 +7874,10 @@ finish_omp_for (location_t locus, enum t
}
if (code == CILK_FOR && i == 0)
orig_decl = decl;
- if (handle_omp_for_class_iterator (i, locus, code, declv, initv,
- condv, incrv, &body, &pre_body,
- clauses, &last, collapse,
- ordered))
+ if (handle_omp_for_class_iterator (i, locus, code, declv, orig_declv,
+ initv, condv, incrv, &body,
+ &pre_body, clauses, &last,
+ collapse, ordered))
return NULL;
continue;
}
@@ -7727,6 +7936,10 @@ finish_omp_for (location_t locus, enum t
omp_for = c_finish_omp_for (locus, code, declv, orig_declv, initv, condv,
incrv, body, pre_body);
+ /* Check for iterators appearing in lb, b or incr expressions. */
+ if (omp_for && !c_omp_check_loop_iv (omp_for, orig_declv, cp_walk_subtrees))
+ omp_for = NULL_TREE;
+
if (omp_for == NULL)
{
if (block)
@@ -7734,6 +7947,8 @@ finish_omp_for (location_t locus, enum t
return NULL;
}
+ add_stmt (omp_for);
+
for (i = 0; i < TREE_VEC_LENGTH (OMP_FOR_INCR (omp_for)); i++)
{
decl = TREE_OPERAND (TREE_VEC_ELT (OMP_FOR_INIT (omp_for), i), 0);
@@ -8013,7 +8228,8 @@ finish_omp_atomic (enum tree_code code,
return;
}
stmt = c_finish_omp_atomic (input_location, code, opcode, lhs, rhs,
- v, lhs1, rhs1, swapped, seq_cst);
+ v, lhs1, rhs1, swapped, seq_cst,
+ processing_template_decl != 0);
if (stmt == error_mark_node)
return;
}
@@ -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,
@@ -4,15 +4,15 @@ void bar (int *);
void
foo (int *p, int q, struct S t, int i, int j, int k, int l)
{
- #pragma omp target map (q), firstprivate (q)
+ #pragma omp target map (q), firstprivate (q) /* { dg-error "appears both in data and map clauses" } */
bar (&q);
#pragma omp target map (p[0]) firstprivate (p) /* { dg-error "appears more than once in data clauses" } */
bar (p);
#pragma omp target firstprivate (p), map (p[0]) /* { dg-error "appears more than once in data clauses" } */
bar (p);
- #pragma omp target map (p[0]) map (p)
+ #pragma omp target map (p[0]) map (p) /* { dg-error "appears both in data and map clauses" } */
bar (p);
- #pragma omp target map (p) , map (p[0])
+ #pragma omp target map (p) , map (p[0]) /* { dg-error "appears both in data and map clauses" } */
bar (p);
#pragma omp target map (q) map (q) /* { dg-error "appears more than once in map clauses" } */
bar (&q);
@@ -24,17 +24,17 @@ foo (int *p, int q, struct S t, int i, i
bar (&t.r);
#pragma omp target map (t.r) map (t.r) /* { dg-error "appears more than once in map clauses" } */
bar (&t.r);
- #pragma omp target firstprivate (t), map (t.r)
+ #pragma omp target firstprivate (t), map (t.r) /* { dg-error "appears both in data and map clauses" } */
bar (&t.r);
- #pragma omp target map (t.r) firstprivate (t)
+ #pragma omp target map (t.r) firstprivate (t) /* { dg-error "appears both in data and map clauses" } */
bar (&t.r);
- #pragma omp target map (t.s[0]) map (t)
+ #pragma omp target map (t.s[0]) map (t) /* { dg-error "appears more than once in map clauses" } */
bar (t.s);
- #pragma omp target map (t) map(t.s[0])
+ #pragma omp target map (t) map(t.s[0]) /* { dg-error "appears more than once in map clauses" } */
bar (t.s);
- #pragma omp target firstprivate (t) map (t.s[0]) /* { dg-error "appears more than once in data clauses" } */
+ #pragma omp target firstprivate (t) map (t.s[0]) /* { dg-error "appears both in data and map clauses" } */
bar (t.s);
- #pragma omp target map (t.s[0]) firstprivate (t) /* { dg-error "appears more than once in data clauses" } */
+ #pragma omp target map (t.s[0]) firstprivate (t) /* { dg-error "appears both in data and map clauses" } */
bar (t.s);
#pragma omp target map (t.s[0]) map (t.s[2]) /* { dg-error "appears more than once in map clauses" } */
bar (t.s);
@@ -46,8 +46,8 @@ foo (int *p, int q, struct S t, int i, i
bar (t.s);
#pragma omp target map (t.r) ,map (t.s[0])
bar (t.s);
- #pragma omp target map (t.r) map (t) map (t.s[0]) firstprivate (t) /* { dg-error "appears more than once in map clauses" } */
- bar (t.s); /* { dg-error "appears more than once in data clauses" "" { target *-*-* } 49 } */
- #pragma omp target map (t) map (t.r) firstprivate (t) map (t.s[0]) /* { dg-error "appears more than once in map clauses" } */
- bar (t.s); /* { dg-error "appears more than once in data clauses" "" { target *-*-* } 51 } */
+ #pragma omp target map (t.r) map (t) map (t.s[0]) firstprivate (t) /* { dg-error "appears both in data and map clauses" } */
+ bar (t.s);
+ #pragma omp target map (t) map (t.r) firstprivate (t) map (t.s[0]) /* { dg-error "appears both in data and map clauses" } */
+ bar (t.s); /* { dg-error "appears more than once in map clauses" "" { target *-*-* } 51 } */
}
@@ -6,7 +6,7 @@ extern int a;
#pragma omp declare target to (a) /* { dg-error "with clauses in between" } */
#pragma omp end declare target
int b;
-#pragma omp declare target to (b) link (b) /* { dg-error "specified both in declare target" } */
+#pragma omp declare target to (b) link (b) /* { dg-error "appears more than once on the same .declare target. directive" } */
int c;
#pragma omp declare target (c)
#pragma omp declare target link (c) /* { dg-error "specified both in declare target" } */
@@ -25,3 +25,21 @@ int g, h;
#pragma omp declare target link (h) /* { dg-error "is threadprivate variable in" } */
int j[10];
#pragma omp declare target to (j[0:4]) /* { dg-error "expected" } */
+int k, l;
+#pragma omp declare target
+int m;
+#pragma omp end declare target
+#pragma omp declare target to (k)
+#pragma omp declare target (k)
+#pragma omp declare target to (k, m) link (l)
+#pragma omp declare target link (l)
+int n, o, s, t;
+#pragma omp declare target to (n) to (n) /* { dg-error "appears more than once on the same .declare target. directive" } */
+#pragma omp declare target link (o, o) /* { dg-error "appears more than once on the same .declare target. directive" } */
+#pragma omp declare target (s, t, s) /* { dg-error "appears more than once on the same .declare target. directive" } */
+int p, q, r;
+#pragma omp declare target (p) to (q) /* { dg-error "expected end of line before .to." } */
+#pragma omp declare target to (p) (q) link (r) /* { dg-error "expected .#pragma omp. clause before" } */
+#pragma omp declare target link (r) (p) /* { dg-error "expected .#pragma omp. clause before" } */
+#pragma omp declare target
+#pragma omp end declare target to (p) /* { dg-error "expected end of line before .to." } */
@@ -355,8 +355,11 @@ test (int n, int o, int p, int q, int r,
int q, i, j;
+#pragma omp declare target
+int s;
+
void
-test2 (int n, int o, int p, int r, int s, int *pp)
+test2 (int n, int o, int p, int r, int *pp)
{
int a[o];
#pragma omp distribute collapse (2) dist_schedule (static, 4) firstprivate (q)
@@ -449,3 +452,4 @@ test2 (int n, int o, int p, int r, int s
s = i * 10;
}
}
+#pragma omp end declare target
@@ -7,14 +7,14 @@ foo (int x)
{
int i = 0;
#pragma omp parallel for simd
- for (i = (i & x); i < 10; i = i + 2)
+ for (i = (i & x); i < 10; i = i + 2) /* { dg-error "initializer expression refers to iteration variable" } */
;
i = 0;
#pragma omp parallel for simd
- for (i = 0; i < (i & x) + 10; i = i + 2)
+ for (i = 0; i < (i & x) + 10; i = i + 2) /* { dg-error "condition expression refers to iteration variable" } */
;
i = 0;
#pragma omp parallel for simd
- for (i = 0; i < 10; i = i + ((i & x) + 2))
+ for (i = 0; i < 10; i = i + ((i & x) + 2)) /* { dg-error "increment expression refers to iteration variable" } */
;
}
@@ -26,9 +26,17 @@ foo (int a[10][10][10], int **b, int x)
bar (a);
#pragma omp parallel reduction(+: a[0:4])
bar (a);
- #pragma omp parallel reduction(+: a[2:4]) /* { dg-error "array section has to be zero-based" } */
+ #pragma omp parallel reduction(+: a[2:4])
bar (a);
- #pragma omp parallel reduction(+: e[2:4]) /* { dg-error "array section has to be zero-based" } */
+ #pragma omp parallel reduction(+: e[2:4])
+ bar (a);
+ #pragma omp parallel reduction(+: a[x:4])
+ bar (a);
+ #pragma omp parallel reduction(+: e[x:4])
+ bar (a);
+ #pragma omp parallel reduction(+: a[x:x])
+ bar (a);
+ #pragma omp parallel reduction(+: e[x:x])
bar (a);
#pragma omp parallel reduction(+: a[0.5:2]) /* { dg-error "low bound \[^\n\r]* of array section does not have integral type" } */
bar (a);
@@ -13,6 +13,13 @@ int f2 (int a, int *b, int c)
return a + *b + c;
}
+/* { dg-final { scan-assembler-times "_ZGVbM8uva32l4_f2:" 1 { target { i?86-*-* x86_64-*-* } } } } */
+/* { dg-final { scan-assembler-times "_ZGVbN8uva32l4_f2:" 1 { target { i?86-*-* x86_64-*-* } } } } */
+/* { dg-final { scan-assembler-times "_ZGVcM8uva32l4_f2:" 1 { target { i?86-*-* x86_64-*-* } } } } */
+/* { dg-final { scan-assembler-times "_ZGVcN8uva32l4_f2:" 1 { target { i?86-*-* x86_64-*-* } } } } */
+/* { dg-final { scan-assembler-times "_ZGVdM8uva32l4_f2:" 1 { target { i?86-*-* x86_64-*-* } } } } */
+/* { dg-final { scan-assembler-times "_ZGVdN8uva32l4_f2:" 1 { target { i?86-*-* x86_64-*-* } } } } */
+
#pragma omp declare simd uniform (a) aligned (b : 8 * sizeof (long long)) linear (c : 4) simdlen (8)
__extension__
long long f3 (long long a, long long *b, long long c);
@@ -40,6 +47,13 @@ f7 (int x)
return x;
}
+/* { dg-final { scan-assembler-times "_ZGVbM16v_f7:" 1 { target { i?86-*-* x86_64-*-* } } } } */
+/* { dg-final { scan-assembler-times "_ZGVbN16v_f7:" 1 { target { i?86-*-* x86_64-*-* } } } } */
+/* { dg-final { scan-assembler-times "_ZGVcM16v_f7:" 1 { target { i?86-*-* x86_64-*-* } } } } */
+/* { dg-final { scan-assembler-times "_ZGVcN16v_f7:" 1 { target { i?86-*-* x86_64-*-* } } } } */
+/* { dg-final { scan-assembler-times "_ZGVdM16v_f7:" 1 { target { i?86-*-* x86_64-*-* } } } } */
+/* { dg-final { scan-assembler-times "_ZGVdN16v_f7:" 1 { target { i?86-*-* x86_64-*-* } } } } */
+
int
f9 (int x)
{
@@ -62,6 +76,13 @@ f13 (int c; int *b; int a; int a, int *b
return a + *b + c;
}
+/* { dg-final { scan-assembler-times "_ZGVbM8uva32l4_f13:" 1 { target { i?86-*-* x86_64-*-* } } } } */
+/* { dg-final { scan-assembler-times "_ZGVbN8uva32l4_f13:" 1 { target { i?86-*-* x86_64-*-* } } } } */
+/* { dg-final { scan-assembler-times "_ZGVcM8uva32l4_f13:" 1 { target { i?86-*-* x86_64-*-* } } } } */
+/* { dg-final { scan-assembler-times "_ZGVcN8uva32l4_f13:" 1 { target { i?86-*-* x86_64-*-* } } } } */
+/* { dg-final { scan-assembler-times "_ZGVdM8uva32l4_f13:" 1 { target { i?86-*-* x86_64-*-* } } } } */
+/* { dg-final { scan-assembler-times "_ZGVdN8uva32l4_f13:" 1 { target { i?86-*-* x86_64-*-* } } } } */
+
#pragma omp declare simd uniform (a) aligned (b : 8 * sizeof (int)) linear (c : 4) simdlen (8)
int
f14 (a, b, c)
@@ -71,6 +92,13 @@ f14 (a, b, c)
return a + *b + c;
}
+/* { dg-final { scan-assembler-times "_ZGVbM8uva32l4_f14:" 1 { target { i?86-*-* x86_64-*-* } } } } */
+/* { dg-final { scan-assembler-times "_ZGVbN8uva32l4_f14:" 1 { target { i?86-*-* x86_64-*-* } } } } */
+/* { dg-final { scan-assembler-times "_ZGVcM8uva32l4_f14:" 1 { target { i?86-*-* x86_64-*-* } } } } */
+/* { dg-final { scan-assembler-times "_ZGVcN8uva32l4_f14:" 1 { target { i?86-*-* x86_64-*-* } } } } */
+/* { dg-final { scan-assembler-times "_ZGVdM8uva32l4_f14:" 1 { target { i?86-*-* x86_64-*-* } } } } */
+/* { dg-final { scan-assembler-times "_ZGVdN8uva32l4_f14:" 1 { target { i?86-*-* x86_64-*-* } } } } */
+
#pragma omp declare simd uniform (a) aligned (b : 8 * sizeof (int)) linear (c : 4) simdlen (8)
int
f15 (int a, int *b, int c)
@@ -78,6 +106,13 @@ f15 (int a, int *b, int c)
return a + *b + c;
}
+/* { dg-final { scan-assembler-times "_ZGVbM8uva32l4_f15:" 1 { target { i?86-*-* x86_64-*-* } } } } */
+/* { dg-final { scan-assembler-times "_ZGVbN8uva32l4_f15:" 1 { target { i?86-*-* x86_64-*-* } } } } */
+/* { dg-final { scan-assembler-times "_ZGVcM8uva32l4_f15:" 1 { target { i?86-*-* x86_64-*-* } } } } */
+/* { dg-final { scan-assembler-times "_ZGVcN8uva32l4_f15:" 1 { target { i?86-*-* x86_64-*-* } } } } */
+/* { dg-final { scan-assembler-times "_ZGVdM8uva32l4_f15:" 1 { target { i?86-*-* x86_64-*-* } } } } */
+/* { dg-final { scan-assembler-times "_ZGVdN8uva32l4_f15:" 1 { target { i?86-*-* x86_64-*-* } } } } */
+
#pragma omp declare simd uniform (d) aligned (e : 8 * sizeof (int)) linear (f : 4) simdlen (8)
int f15 (int d, int *e, int f);
@@ -90,6 +125,19 @@ int f17 (int g, long *h)
return g + h[0];
}
+/* { dg-final { scan-assembler-times "_ZGVbM4l20va8_f17:" 1 { target { { i?86-*-* x86_64-*-* } && lp64 } } } } */
+/* { dg-final { scan-assembler-times "_ZGVbN4l20va8_f17:" 1 { target { { i?86-*-* x86_64-*-* } && lp64 } } } } */
+/* { dg-final { scan-assembler-times "_ZGVcM4l20va8_f17:" 1 { target { { i?86-*-* x86_64-*-* } && lp64 } } } } */
+/* { dg-final { scan-assembler-times "_ZGVcN4l20va8_f17:" 1 { target { { i?86-*-* x86_64-*-* } && lp64 } } } } */
+/* { dg-final { scan-assembler-times "_ZGVdM4l20va8_f17:" 1 { target { { i?86-*-* x86_64-*-* } && lp64 } } } } */
+/* { dg-final { scan-assembler-times "_ZGVdN4l20va8_f17:" 1 { target { { i?86-*-* x86_64-*-* } && lp64 } } } } */
+/* { dg-final { scan-assembler-times "_ZGVbM4l12va4_f17:" 1 { target { { i?86-*-* x86_64-*-* } && ilp32 } } } } */
+/* { dg-final { scan-assembler-times "_ZGVbN4l12va4_f17:" 1 { target { { i?86-*-* x86_64-*-* } && ilp32 } } } } */
+/* { dg-final { scan-assembler-times "_ZGVcM4l12va4_f17:" 1 { target { { i?86-*-* x86_64-*-* } && ilp32 } } } } */
+/* { dg-final { scan-assembler-times "_ZGVcN4l12va4_f17:" 1 { target { { i?86-*-* x86_64-*-* } && ilp32 } } } } */
+/* { dg-final { scan-assembler-times "_ZGVdM4l12va4_f17:" 1 { target { { i?86-*-* x86_64-*-* } && ilp32 } } } } */
+/* { dg-final { scan-assembler-times "_ZGVdN4l12va4_f17:" 1 { target { { i?86-*-* x86_64-*-* } && ilp32 } } } } */
+
#pragma omp declare simd aligned (i : sizeof (*i)) linear (j : 2 * sizeof (i[0]) + sizeof (j)) simdlen (4)
int
f18 (j, i)
@@ -98,3 +146,16 @@ f18 (j, i)
{
return j + i[0];
}
+
+/* { dg-final { scan-assembler-times "_ZGVbM4l20va8_f18:" 1 { target { { i?86-*-* x86_64-*-* } && lp64 } } } } */
+/* { dg-final { scan-assembler-times "_ZGVbN4l20va8_f18:" 1 { target { { i?86-*-* x86_64-*-* } && lp64 } } } } */
+/* { dg-final { scan-assembler-times "_ZGVcM4l20va8_f18:" 1 { target { { i?86-*-* x86_64-*-* } && lp64 } } } } */
+/* { dg-final { scan-assembler-times "_ZGVcN4l20va8_f18:" 1 { target { { i?86-*-* x86_64-*-* } && lp64 } } } } */
+/* { dg-final { scan-assembler-times "_ZGVdM4l20va8_f18:" 1 { target { { i?86-*-* x86_64-*-* } && lp64 } } } } */
+/* { dg-final { scan-assembler-times "_ZGVdN4l20va8_f18:" 1 { target { { i?86-*-* x86_64-*-* } && lp64 } } } } */
+/* { dg-final { scan-assembler-times "_ZGVbM4l12va4_f18:" 1 { target { { i?86-*-* x86_64-*-* } && ilp32 } } } } */
+/* { dg-final { scan-assembler-times "_ZGVbN4l12va4_f18:" 1 { target { { i?86-*-* x86_64-*-* } && ilp32 } } } } */
+/* { dg-final { scan-assembler-times "_ZGVcM4l12va4_f18:" 1 { target { { i?86-*-* x86_64-*-* } && ilp32 } } } } */
+/* { dg-final { scan-assembler-times "_ZGVcN4l12va4_f18:" 1 { target { { i?86-*-* x86_64-*-* } && ilp32 } } } } */
+/* { dg-final { scan-assembler-times "_ZGVdM4l12va4_f18:" 1 { target { { i?86-*-* x86_64-*-* } && ilp32 } } } } */
+/* { dg-final { scan-assembler-times "_ZGVdN4l12va4_f18:" 1 { target { { i?86-*-* x86_64-*-* } && ilp32 } } } } */
@@ -12,7 +12,8 @@ void T::test()
#pragma omp parallel private(n)
n = 1;
- #pragma omp parallel shared(n) // { dg-error "T::n" }
+ #pragma omp parallel shared(n)
+ #pragma omp single
n = 1;
#pragma omp parallel firstprivate(n)
@@ -1,5 +1,6 @@
// Test parsing of #pragma omp declare simd
// { dg-do compile }
+// { dg-options "-fopenmp -ffat-lto-objects" }
#pragma omp declare simd uniform (a) aligned (b : 8 * sizeof (int)) \
linear (c : 4) simdlen (8) notinbranch
@@ -13,6 +14,13 @@ int f2 (int a, int *b, int c)
return a + *b + c;
}
+// { dg-final { scan-assembler-times "_ZGVbM8uva32l4__Z2f2iPii:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVbN8uva32l4__Z2f2iPii:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVcM8uva32l4__Z2f2iPii:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVcN8uva32l4__Z2f2iPii:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVdM8uva32l4__Z2f2iPii:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVdN8uva32l4__Z2f2iPii:" 1 { target { i?86-*-* x86_64-*-* } } } }
+
#pragma omp declare simd uniform (c) aligned (b : 4 * sizeof (int)) linear (a : 4) simdlen (4)
template <typename T>
T f3 (int a, int *b, T c);
@@ -71,6 +79,13 @@ namespace N1
}
}
+// { dg-final { scan-assembler-times "_ZGVbM2va16__ZN2N12N23f10EPx:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVbN2va16__ZN2N12N23f10EPx:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVcM2va16__ZN2N12N23f10EPx:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVcN2va16__ZN2N12N23f10EPx:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVdM2va16__ZN2N12N23f10EPx:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVdN2va16__ZN2N12N23f10EPx:" 1 { target { i?86-*-* x86_64-*-* } } } }
+
struct A
{
#pragma omp declare simd uniform (a) aligned (b : 8 * sizeof (int)) linear (c : 4) simdlen (8)
@@ -172,6 +187,13 @@ int B<int>::f25<7> (int a, int *b, int c
return a + *b + c;
}
+// { dg-final { scan-assembler-times "_ZGVbM8vuva32u__ZN1BIiE3f25ILi7EEEiiPii:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVbN8vuva32u__ZN1BIiE3f25ILi7EEEiiPii:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVcM8vuva32u__ZN1BIiE3f25ILi7EEEiiPii:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVcN8vuva32u__ZN1BIiE3f25ILi7EEEiiPii:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVdM8vuva32u__ZN1BIiE3f25ILi7EEEiiPii:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVdN8vuva32u__ZN1BIiE3f25ILi7EEEiiPii:" 1 { target { i?86-*-* x86_64-*-* } } } }
+
#pragma omp declare simd simdlen (4) aligned (b : 8 * sizeof (int)) linear (a, c : 2)
template <>
template <>
@@ -180,6 +202,13 @@ int B<int>::f26<-1> (int a, int *b, int
return a + *b + c;
}
+// { dg-final { scan-assembler-times "_ZGVbM4vl2va32__ZN1BIiE3f26ILin1EEEiiPii:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVbN4vl2va32__ZN1BIiE3f26ILin1EEEiiPii:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVcM4vl2va32__ZN1BIiE3f26ILin1EEEiiPii:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVcN4vl2va32__ZN1BIiE3f26ILin1EEEiiPii:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVdM4vl2va32__ZN1BIiE3f26ILin1EEEiiPii:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVdN4vl2va32__ZN1BIiE3f26ILin1EEEiiPii:" 1 { target { i?86-*-* x86_64-*-* } } } }
+
int
f27 (int x)
{
@@ -202,6 +231,13 @@ f30 (int x)
return x;
}
+// { dg-final { scan-assembler-times "_ZGVbM16v__Z3f30i:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVbN16v__Z3f30i:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVcM16v__Z3f30i:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVcN16v__Z3f30i:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVdM16v__Z3f30i:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVdN16v__Z3f30i:" 1 { target { i?86-*-* x86_64-*-* } } } }
+
template <int N>
struct C
{
@@ -1,4 +1,5 @@
// { dg-do compile }
+// { dg-options "-fopenmp -ffat-lto-objects" }
#pragma omp declare simd uniform(b) linear(c, d) linear(uval(e)) linear(ref(f))
int f1 (int a, int b, int c, int &d, int &e, int &f)
@@ -12,6 +13,13 @@ int f1 (int a, int b, int c, int &d, int
return a + b + c + d + e + f;
}
+// { dg-final { scan-assembler-times "_ZGVbM4vulLUR4__Z2f1iiiRiS_S_:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVbN4vulLUR4__Z2f1iiiRiS_S_:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVcM4vulLUR4__Z2f1iiiRiS_S_:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVcN4vulLUR4__Z2f1iiiRiS_S_:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVdM8vulLUR4__Z2f1iiiRiS_S_:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVdN8vulLUR4__Z2f1iiiRiS_S_:" 1 { target { i?86-*-* x86_64-*-* } } } }
+
#pragma omp declare simd uniform(b) linear(c, d) linear(uval(e)) linear(ref(f))
int f2 (int a, int b, int c, int &d, int &e, int &f)
{
@@ -30,12 +38,26 @@ int f2 (int a, int b, int c, int &d, int
return a + b + c + d + e + f;
}
+// { dg-final { scan-assembler-times "_ZGVbM4vulLUR4__Z2f2iiiRiS_S_:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVbN4vulLUR4__Z2f2iiiRiS_S_:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVcM4vulLUR4__Z2f2iiiRiS_S_:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVcN4vulLUR4__Z2f2iiiRiS_S_:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVdM8vulLUR4__Z2f2iiiRiS_S_:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVdN8vulLUR4__Z2f2iiiRiS_S_:" 1 { target { i?86-*-* x86_64-*-* } } } }
+
#pragma omp declare simd uniform(b) linear(c, d) linear(uval(e)) linear(ref(f))
int f3 (const int a, const int b, const int c, const int &d, const int &e, const int &f)
{
return a + b + c + d + e + f;
}
+// { dg-final { scan-assembler-times "_ZGVbM4vulLUR4__Z2f3iiiRKiS0_S0_:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVbN4vulLUR4__Z2f3iiiRKiS0_S0_:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVcM4vulLUR4__Z2f3iiiRKiS0_S0_:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVcN4vulLUR4__Z2f3iiiRKiS0_S0_:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVdM8vulLUR4__Z2f3iiiRKiS0_S0_:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVdN8vulLUR4__Z2f3iiiRKiS0_S0_:" 1 { target { i?86-*-* x86_64-*-* } } } }
+
#pragma omp declare simd uniform(b) linear(c, d) linear(uval(e)) linear(ref(f))
int f4 (const int a, const int b, const int c, const int &d, const int &e, const int &f)
{
@@ -47,3 +69,10 @@ int f4 (const int a, const int b, const
asm volatile ("" : : "r" (&f));
return a + b + c + d + e + f;
}
+
+// { dg-final { scan-assembler-times "_ZGVbM4vulLUR4__Z2f4iiiRKiS0_S0_:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVbN4vulLUR4__Z2f4iiiRKiS0_S0_:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVcM4vulLUR4__Z2f4iiiRKiS0_S0_:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVcN4vulLUR4__Z2f4iiiRKiS0_S0_:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVdM8vulLUR4__Z2f4iiiRKiS0_S0_:" 1 { target { i?86-*-* x86_64-*-* } } } }
+// { dg-final { scan-assembler-times "_ZGVdN8vulLUR4__Z2f4iiiRKiS0_S0_:" 1 { target { i?86-*-* x86_64-*-* } } } }
@@ -76,19 +76,19 @@ B::m2 ()
#pragma omp parallel for reduction (+:g) // { dg-error "has const type for .reduction." }
for (int i = 0; i < 10; i++)
;
- #pragma omp parallel shared (a) // { dg-error "is not a variable in clause" }
+ #pragma omp parallel shared (a)
;
- #pragma omp parallel shared (b) // { dg-error "is not a variable in clause" }
+ #pragma omp parallel shared (b)
;
- #pragma omp parallel shared (c) // { dg-error "is not a variable in clause" }
+ #pragma omp parallel shared (c)
;
- #pragma omp parallel shared (e) // { dg-error "is not a variable in clause" }
+ #pragma omp parallel shared (e)
;
- #pragma omp parallel shared (f) // { dg-error "is not a variable in clause" }
+ #pragma omp parallel shared (f)
;
- #pragma omp parallel shared (g) // { dg-error "is not a variable in clause" }
+ #pragma omp parallel shared (g)
;
- #pragma omp parallel shared (h) // { dg-error "is not a variable in clause" }
+ #pragma omp parallel shared (h) // { dg-error "is predetermined .shared. for .shared." }
;
return 0;
}
@@ -150,19 +150,19 @@ B::m4 () const
#pragma omp parallel for reduction (+:g) // { dg-error "has const type for .reduction." }
for (int i = 0; i < 10; i++)
;
- #pragma omp parallel shared (a) // { dg-error "is not a variable in clause" }
+ #pragma omp parallel shared (a) // { dg-error "is predetermined .shared. for .shared." }
;
- #pragma omp parallel shared (b) // { dg-error "is not a variable in clause" }
+ #pragma omp parallel shared (b)
;
- #pragma omp parallel shared (c) // { dg-error "is not a variable in clause" }
+ #pragma omp parallel shared (c)
;
- #pragma omp parallel shared (e) // { dg-error "is not a variable in clause" }
+ #pragma omp parallel shared (e)
;
- #pragma omp parallel shared (f) // { dg-error "is not a variable in clause" }
+ #pragma omp parallel shared (f)
;
- #pragma omp parallel shared (g) // { dg-error "is not a variable in clause" }
+ #pragma omp parallel shared (g)
;
- #pragma omp parallel shared (h) // { dg-error "is not a variable in clause" }
+ #pragma omp parallel shared (h) // { dg-error "is predetermined .shared. for .shared." }
;
return 0;
}
@@ -6,12 +6,12 @@ template<int> void
foo ()
{
#pragma omp parallel for
- for (auto i = i = 0; i<4; ++i) // { dg-error "incomplete|unable|invalid|auto" }
+ for (auto i = i = 0; i<4; ++i) // { dg-error "initializer expression refers to iteration variable" }
;
}
void
bar ()
{
- foo<0> (); // { dg-message "required from here" }
+ foo<0> ();
}
@@ -39,10 +39,9 @@
/* Special map kinds, enumerated starting here. */
#define GOMP_MAP_FLAG_SPECIAL_0 (1 << 2)
#define GOMP_MAP_FLAG_SPECIAL_1 (1 << 3)
+#define GOMP_MAP_FLAG_SPECIAL_2 (1 << 4)
#define GOMP_MAP_FLAG_SPECIAL (GOMP_MAP_FLAG_SPECIAL_1 \
| GOMP_MAP_FLAG_SPECIAL_0)
-/* OpenMP always flag. */
-#define GOMP_MAP_FLAG_ALWAYS (1 << 6)
/* Flag to force a specific behavior (or else, trigger a run-time error). */
#define GOMP_MAP_FLAG_FORCE (1 << 7)
@@ -95,34 +94,43 @@ enum gomp_map_kind
GOMP_MAP_FORCE_TOFROM = (GOMP_MAP_FLAG_FORCE | GOMP_MAP_TOFROM),
/* If not already present, allocate. And unconditionally copy to
device. */
- GOMP_MAP_ALWAYS_TO = (GOMP_MAP_FLAG_ALWAYS | GOMP_MAP_TO),
+ GOMP_MAP_ALWAYS_TO = (GOMP_MAP_FLAG_SPECIAL_2 | GOMP_MAP_TO),
/* If not already present, allocate. And unconditionally copy from
device. */
- GOMP_MAP_ALWAYS_FROM = (GOMP_MAP_FLAG_ALWAYS | GOMP_MAP_FROM),
+ GOMP_MAP_ALWAYS_FROM = (GOMP_MAP_FLAG_SPECIAL_2
+ | GOMP_MAP_FROM),
/* If not already present, allocate. And unconditionally copy to and from
device. */
- GOMP_MAP_ALWAYS_TOFROM = (GOMP_MAP_FLAG_ALWAYS | GOMP_MAP_TOFROM),
+ GOMP_MAP_ALWAYS_TOFROM = (GOMP_MAP_FLAG_SPECIAL_2
+ | GOMP_MAP_TOFROM),
/* Map a sparse struct; the address is the base of the structure, alignment
it's required alignment, and size is the number of adjacent entries
that belong to the struct. The adjacent entries should be sorted by
increasing address, so it is easy to determine lowest needed address
(address of the first adjacent entry) and highest needed address
(address of the last adjacent entry plus its size). */
- GOMP_MAP_STRUCT = (GOMP_MAP_FLAG_ALWAYS
+ GOMP_MAP_STRUCT = (GOMP_MAP_FLAG_SPECIAL_2
| GOMP_MAP_FLAG_SPECIAL | 0),
+ /* On a location of a pointer/reference that is assumed to be already mapped
+ earlier, store the translated address of the preceeding mapping.
+ No refcount is bumped by this, and the store is done unconditionally. */
+ GOMP_MAP_ALWAYS_POINTER = (GOMP_MAP_FLAG_SPECIAL_2
+ | GOMP_MAP_FLAG_SPECIAL | 1),
/* Forced deallocation of zero length array section. */
GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION
- = (GOMP_MAP_FLAG_ALWAYS
+ = (GOMP_MAP_FLAG_SPECIAL_2
| GOMP_MAP_FLAG_SPECIAL | 3),
- /* OpenMP 4.1 alias for forced deallocation. */
+ /* OpenMP 4.5 alias for forced deallocation. */
GOMP_MAP_DELETE = GOMP_MAP_FORCE_DEALLOC,
/* Decrement usage count and deallocate if zero. */
- GOMP_MAP_RELEASE = (GOMP_MAP_FLAG_ALWAYS
+ GOMP_MAP_RELEASE = (GOMP_MAP_FLAG_SPECIAL_2
| GOMP_MAP_FORCE_DEALLOC),
/* Internal to GCC, not used in libgomp. */
/* Do not map, but pointer assign a pointer instead. */
- GOMP_MAP_FIRSTPRIVATE_POINTER = (GOMP_MAP_LAST | 1)
+ GOMP_MAP_FIRSTPRIVATE_POINTER = (GOMP_MAP_LAST | 1),
+ /* Do not map, but pointer assign a reference instead. */
+ GOMP_MAP_FIRSTPRIVATE_REFERENCE = (GOMP_MAP_LAST | 2)
};
#define GOMP_MAP_COPY_TO_P(X) \
@@ -142,6 +150,9 @@ enum gomp_map_kind
#define GOMP_MAP_ALWAYS_FROM_P(X) \
(((X) == GOMP_MAP_ALWAYS_FROM) || ((X) == GOMP_MAP_ALWAYS_TOFROM))
+#define GOMP_MAP_ALWAYS_P(X) \
+ (GOMP_MAP_ALWAYS_TO_P (X) || ((X) == GOMP_MAP_ALWAYS_FROM))
+
/* Asynchronous behavior. Keep in sync with
libgomp/{openacc.h,openacc.f90,openacc_lib.h}:acc_async_t. */
@@ -175,7 +186,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. */
@@ -52,6 +52,10 @@ extern bool GOMP_loop_static_start (long
extern bool GOMP_loop_dynamic_start (long, long, long, long, long *, long *);
extern bool GOMP_loop_guided_start (long, long, long, long, long *, long *);
extern bool GOMP_loop_runtime_start (long, long, long, long *, long *);
+extern bool GOMP_loop_nonmonotonic_dynamic_start (long, long, long, long,
+ long *, long *);
+extern bool GOMP_loop_nonmonotonic_guided_start (long, long, long, long,
+ long *, long *);
extern bool GOMP_loop_ordered_static_start (long, long, long, long,
long *, long *);
@@ -65,6 +69,8 @@ extern bool GOMP_loop_static_next (long
extern bool GOMP_loop_dynamic_next (long *, long *);
extern bool GOMP_loop_guided_next (long *, long *);
extern bool GOMP_loop_runtime_next (long *, long *);
+extern bool GOMP_loop_nonmonotonic_dynamic_next (long *, long *);
+extern bool GOMP_loop_nonmonotonic_guided_next (long *, long *);
extern bool GOMP_loop_ordered_static_next (long *, long *);
extern bool GOMP_loop_ordered_dynamic_next (long *, long *);
@@ -100,6 +106,12 @@ extern void GOMP_parallel_loop_guided (v
extern void GOMP_parallel_loop_runtime (void (*)(void *), void *,
unsigned, long, long, long,
unsigned);
+extern void GOMP_parallel_loop_nonmonotonic_dynamic (void (*)(void *), void *,
+ unsigned, long, long,
+ long, long, unsigned);
+extern void GOMP_parallel_loop_nonmonotonic_guided (void (*)(void *), void *,
+ unsigned, long, long,
+ long, long, unsigned);
extern void GOMP_loop_end (void);
extern void GOMP_loop_end_nowait (void);
@@ -130,6 +142,18 @@ extern bool GOMP_loop_ull_runtime_start
unsigned long long,
unsigned long long *,
unsigned long long *);
+extern bool GOMP_loop_ull_nonmonotonic_dynamic_start (bool, unsigned long long,
+ unsigned long long,
+ unsigned long long,
+ unsigned long long,
+ unsigned long long *,
+ unsigned long long *);
+extern bool GOMP_loop_ull_nonmonotonic_guided_start (bool, unsigned long long,
+ unsigned long long,
+ unsigned long long,
+ unsigned long long,
+ unsigned long long *,
+ unsigned long long *);
extern bool GOMP_loop_ull_ordered_static_start (bool, unsigned long long,
unsigned long long,
@@ -163,6 +187,10 @@ extern bool GOMP_loop_ull_guided_next (u
unsigned long long *);
extern bool GOMP_loop_ull_runtime_next (unsigned long long *,
unsigned long long *);
+extern bool GOMP_loop_ull_nonmonotonic_dynamic_next (unsigned long long *,
+ unsigned long long *);
+extern bool GOMP_loop_ull_nonmonotonic_guided_next (unsigned long long *,
+ unsigned long long *);
extern bool GOMP_loop_ull_ordered_static_next (unsigned long long *,
unsigned long long *);
@@ -249,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 **);
@@ -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;
@@ -286,6 +286,16 @@ GOMP_4.5 {
GOMP_loop_ull_doacross_static_start;
GOMP_doacross_ull_post;
GOMP_doacross_ull_wait;
+ GOMP_loop_nonmonotonic_dynamic_next;
+ GOMP_loop_nonmonotonic_dynamic_start;
+ GOMP_loop_nonmonotonic_guided_next;
+ GOMP_loop_nonmonotonic_guided_start;
+ GOMP_loop_ull_nonmonotonic_dynamic_next;
+ GOMP_loop_ull_nonmonotonic_dynamic_start;
+ GOMP_loop_ull_nonmonotonic_guided_next;
+ GOMP_loop_ull_nonmonotonic_guided_start;
+ GOMP_parallel_loop_nonmonotonic_dynamic;
+ GOMP_parallel_loop_nonmonotonic_guided;
} GOMP_4.0.1;
OACC_2.0 {
@@ -111,6 +111,11 @@ gomp_loop_static_start (long start, long
return !gomp_iter_static_next (istart, iend);
}
+/* The current dynamic implementation is always monotonic. The
+ entrypoints without nonmonotonic in them have to be always monotonic,
+ but the nonmonotonic ones could be changed to use work-stealing for
+ improved scalability. */
+
static bool
gomp_loop_dynamic_start (long start, long end, long incr, long chunk_size,
long *istart, long *iend)
@@ -136,6 +141,9 @@ gomp_loop_dynamic_start (long start, lon
return ret;
}
+/* Similarly as for dynamic, though the question is how can the chunk sizes
+ be decreased without a central locking or atomics. */
+
static bool
gomp_loop_guided_start (long start, long end, long incr, long chunk_size,
long *istart, long *iend)
@@ -630,6 +638,37 @@ GOMP_parallel_loop_guided (void (*fn) (v
GOMP_parallel_end ();
}
+#ifdef HAVE_ATTRIBUTE_ALIAS
+extern __typeof(GOMP_parallel_loop_dynamic) GOMP_parallel_loop_nonmonotonic_dynamic
+ __attribute__((alias ("GOMP_parallel_loop_dynamic")));
+extern __typeof(GOMP_parallel_loop_guided) GOMP_parallel_loop_nonmonotonic_guided
+ __attribute__((alias ("GOMP_parallel_loop_guided")));
+#else
+void
+GOMP_parallel_loop_nonmonotonic_dynamic (void (*fn) (void *), void *data,
+ unsigned num_threads, long start,
+ long end, long incr, long chunk_size,
+ unsigned flags)
+{
+ gomp_parallel_loop_start (fn, data, num_threads, start, end, incr,
+ GFS_DYNAMIC, chunk_size, flags);
+ fn (data);
+ GOMP_parallel_end ();
+}
+
+void
+GOMP_parallel_loop_nonmonotonic_guided (void (*fn) (void *), void *data,
+ unsigned num_threads, long start,
+ long end, long incr, long chunk_size,
+ unsigned flags)
+{
+ gomp_parallel_loop_start (fn, data, num_threads, start, end, incr,
+ GFS_GUIDED, chunk_size, flags);
+ fn (data);
+ GOMP_parallel_end ();
+}
+#endif
+
void
GOMP_parallel_loop_runtime (void (*fn) (void *), void *data,
unsigned num_threads, long start, long end,
@@ -678,6 +717,10 @@ extern __typeof(gomp_loop_dynamic_start)
__attribute__((alias ("gomp_loop_dynamic_start")));
extern __typeof(gomp_loop_guided_start) GOMP_loop_guided_start
__attribute__((alias ("gomp_loop_guided_start")));
+extern __typeof(gomp_loop_dynamic_start) GOMP_loop_nonmonotonic_dynamic_start
+ __attribute__((alias ("gomp_loop_dynamic_start")));
+extern __typeof(gomp_loop_guided_start) GOMP_loop_nonmonotonic_guided_start
+ __attribute__((alias ("gomp_loop_guided_start")));
extern __typeof(gomp_loop_ordered_static_start) GOMP_loop_ordered_static_start
__attribute__((alias ("gomp_loop_ordered_static_start")));
@@ -699,6 +742,10 @@ extern __typeof(gomp_loop_dynamic_next)
__attribute__((alias ("gomp_loop_dynamic_next")));
extern __typeof(gomp_loop_guided_next) GOMP_loop_guided_next
__attribute__((alias ("gomp_loop_guided_next")));
+extern __typeof(gomp_loop_dynamic_next) GOMP_loop_nonmonotonic_dynamic_next
+ __attribute__((alias ("gomp_loop_dynamic_next")));
+extern __typeof(gomp_loop_guided_next) GOMP_loop_nonmonotonic_guided_next
+ __attribute__((alias ("gomp_loop_guided_next")));
extern __typeof(gomp_loop_ordered_static_next) GOMP_loop_ordered_static_next
__attribute__((alias ("gomp_loop_ordered_static_next")));
@@ -729,6 +776,21 @@ GOMP_loop_guided_start (long start, long
}
bool
+GOMP_loop_nonmonotonic_dynamic_start (long start, long end, long incr,
+ long chunk_size, long *istart,
+ long *iend)
+{
+ return gomp_loop_dynamic_start (start, end, incr, chunk_size, istart, iend);
+}
+
+bool
+GOMP_loop_nonmonotonic_guided_start (long start, long end, long incr,
+ long chunk_size, long *istart, long *iend)
+{
+ return gomp_loop_guided_start (start, end, incr, chunk_size, istart, iend);
+}
+
+bool
GOMP_loop_ordered_static_start (long start, long end, long incr,
long chunk_size, long *istart, long *iend)
{
@@ -793,6 +855,18 @@ GOMP_loop_guided_next (long *istart, lon
{
return gomp_loop_guided_next (istart, iend);
}
+
+bool
+GOMP_loop_nonmonotonic_dynamic_next (long *istart, long *iend)
+{
+ return gomp_loop_dynamic_next (istart, iend);
+}
+
+bool
+GOMP_loop_nonmonotonic_guided_next (long *istart, long *iend)
+{
+ return gomp_loop_guided_next (istart, iend);
+}
bool
GOMP_loop_ordered_static_next (long *istart, long *iend)
@@ -566,6 +566,10 @@ extern __typeof(gomp_loop_ull_dynamic_st
__attribute__((alias ("gomp_loop_ull_dynamic_start")));
extern __typeof(gomp_loop_ull_guided_start) GOMP_loop_ull_guided_start
__attribute__((alias ("gomp_loop_ull_guided_start")));
+extern __typeof(gomp_loop_ull_dynamic_start) GOMP_loop_ull_nonmonotonic_dynamic_start
+ __attribute__((alias ("gomp_loop_ull_dynamic_start")));
+extern __typeof(gomp_loop_ull_guided_start) GOMP_loop_ull_nonmonotonic_guided_start
+ __attribute__((alias ("gomp_loop_ull_guided_start")));
extern __typeof(gomp_loop_ull_ordered_static_start) GOMP_loop_ull_ordered_static_start
__attribute__((alias ("gomp_loop_ull_ordered_static_start")));
@@ -587,6 +591,10 @@ extern __typeof(gomp_loop_ull_dynamic_ne
__attribute__((alias ("gomp_loop_ull_dynamic_next")));
extern __typeof(gomp_loop_ull_guided_next) GOMP_loop_ull_guided_next
__attribute__((alias ("gomp_loop_ull_guided_next")));
+extern __typeof(gomp_loop_ull_dynamic_next) GOMP_loop_ull_nonmonotonic_dynamic_next
+ __attribute__((alias ("gomp_loop_ull_dynamic_next")));
+extern __typeof(gomp_loop_ull_guided_next) GOMP_loop_ull_nonmonotonic_guided_next
+ __attribute__((alias ("gomp_loop_ull_guided_next")));
extern __typeof(gomp_loop_ull_ordered_static_next) GOMP_loop_ull_ordered_static_next
__attribute__((alias ("gomp_loop_ull_ordered_static_next")));
@@ -623,6 +631,25 @@ GOMP_loop_ull_guided_start (bool up, gom
}
bool
+GOMP_loop_ull_nonmonotonic_dynamic_start (bool up, gomp_ull start,
+ gomp_ull end, gomp_ull incr,
+ gomp_ull chunk_size,
+ gomp_ull *istart, gomp_ull *iend)
+{
+ return gomp_loop_ull_dynamic_start (up, start, end, incr, chunk_size, istart,
+ iend);
+}
+
+bool
+GOMP_loop_ull_nonmonotonic_guided_start (bool up, gomp_ull start, gomp_ull end,
+ gomp_ull incr, gomp_ull chunk_size,
+ gomp_ull *istart, gomp_ull *iend)
+{
+ return gomp_loop_ull_guided_start (up, start, end, incr, chunk_size, istart,
+ iend);
+}
+
+bool
GOMP_loop_ull_ordered_static_start (bool up, gomp_ull start, gomp_ull end,
gomp_ull incr, gomp_ull chunk_size,
gomp_ull *istart, gomp_ull *iend)
@@ -693,6 +720,18 @@ GOMP_loop_ull_guided_next (gomp_ull *ist
{
return gomp_loop_ull_guided_next (istart, iend);
}
+
+bool
+GOMP_loop_ull_nonmonotonic_dynamic_next (gomp_ull *istart, gomp_ull *iend)
+{
+ return gomp_loop_ull_dynamic_next (istart, iend);
+}
+
+bool
+GOMP_loop_ull_nonmonotonic_guided_next (gomp_ull *istart, gomp_ull *iend)
+{
+ return gomp_loop_ull_guided_next (istart, iend);
+}
bool
GOMP_loop_ull_ordered_static_next (gomp_ull *istart, gomp_ull *iend)
@@ -162,7 +162,20 @@ gomp_map_lookup (splay_tree mem_map, spl
return splay_tree_lookup (mem_map, key);
}
-/* Handle the case where gomp_map_lookup found oldn for newn.
+static inline splay_tree_key
+gomp_map_0len_lookup (splay_tree mem_map, splay_tree_key key)
+{
+ if (key->host_start != key->host_end)
+ return splay_tree_lookup (mem_map, key);
+
+ key->host_end++;
+ splay_tree_key n = splay_tree_lookup (mem_map, key);
+ key->host_end--;
+ return n;
+}
+
+/* Handle the case where gomp_map_lookup, splay_tree_lookup or
+ gomp_map_0len_lookup found oldn for newn.
Helper function of gomp_map_vars. */
static inline void
@@ -306,6 +319,26 @@ gomp_map_fields_existing (struct target_
(void *) cur_node.host_end);
}
+static inline uintptr_t
+gomp_map_val (struct target_mem_desc *tgt, void **hostaddrs, size_t i)
+{
+ if (tgt->list[i].key != NULL)
+ return tgt->list[i].key->tgt->tgt_start
+ + tgt->list[i].key->tgt_offset
+ + tgt->list[i].offset;
+ if (tgt->list[i].offset == ~(uintptr_t) 0)
+ return (uintptr_t) hostaddrs[i];
+ if (tgt->list[i].offset == ~(uintptr_t) 1)
+ return 0;
+ if (tgt->list[i].offset == ~(uintptr_t) 2)
+ return tgt->list[i + 1].key->tgt->tgt_start
+ + tgt->list[i + 1].key->tgt_offset
+ + tgt->list[i + 1].offset
+ + (uintptr_t) hostaddrs[i]
+ - (uintptr_t) hostaddrs[i + 1];
+ return tgt->tgt_start + tgt->list[i].offset;
+}
+
attribute_hidden struct target_mem_desc *
gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
void **hostaddrs, void **devaddrs, size_t *sizes, void *kinds,
@@ -400,6 +433,13 @@ gomp_map_vars (struct gomp_device_descr
i--;
continue;
}
+ else if ((kind & typemask) == GOMP_MAP_ALWAYS_POINTER)
+ {
+ tgt->list[i].key = NULL;
+ tgt->list[i].offset = ~(uintptr_t) 1;
+ has_firstprivate = true;
+ continue;
+ }
cur_node.host_start = (uintptr_t) hostaddrs[i];
if (!GOMP_MAP_POINTER_P (kind & typemask))
cur_node.host_end = cur_node.host_start + sizes[i];
@@ -420,7 +460,7 @@ gomp_map_vars (struct gomp_device_descr
splay_tree_key n;
if ((kind & typemask) == GOMP_MAP_ZERO_LEN_ARRAY_SECTION)
{
- n = gomp_map_lookup (mem_map, &cur_node);
+ n = gomp_map_0len_lookup (mem_map, &cur_node);
if (!n)
{
tgt->list[i].key = NULL;
@@ -558,6 +598,32 @@ gomp_map_vars (struct gomp_device_descr
sizes, kinds);
i--;
continue;
+ case GOMP_MAP_ALWAYS_POINTER:
+ cur_node.host_start = (uintptr_t) hostaddrs[i];
+ cur_node.host_end = cur_node.host_start + sizeof (void *);
+ n = splay_tree_lookup (mem_map, &cur_node);
+ if (n == NULL
+ || n->host_start > cur_node.host_start
+ || n->host_end < cur_node.host_end)
+ {
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_fatal ("always pointer not mapped");
+ }
+ if ((get_kind (short_mapkind, kinds, i - 1) & typemask)
+ != GOMP_MAP_ALWAYS_POINTER)
+ cur_node.tgt_offset = gomp_map_val (tgt, hostaddrs, i - 1);
+ if (cur_node.tgt_offset)
+ cur_node.tgt_offset -= sizes[i];
+ devicep->host2dev_func (devicep->target_id,
+ (void *) (n->tgt->tgt_start
+ + n->tgt_offset
+ + cur_node.host_start
+ - n->host_start),
+ (void *) &cur_node.tgt_offset,
+ sizeof (void *));
+ cur_node.tgt_offset = n->tgt->tgt_start + n->tgt_offset
+ + cur_node.host_start - n->host_start;
+ continue;
default:
break;
}
@@ -701,26 +767,7 @@ gomp_map_vars (struct gomp_device_descr
{
for (i = 0; i < mapnum; i++)
{
- if (tgt->list[i].key == NULL)
- {
- if (tgt->list[i].offset == ~(uintptr_t) 0)
- cur_node.tgt_offset = (uintptr_t) hostaddrs[i];
- else if (tgt->list[i].offset == ~(uintptr_t) 1)
- cur_node.tgt_offset = 0;
- else if (tgt->list[i].offset == ~(uintptr_t) 2)
- cur_node.tgt_offset = tgt->list[i + 1].key->tgt->tgt_start
- + tgt->list[i + 1].key->tgt_offset
- + tgt->list[i + 1].offset
- + (uintptr_t) hostaddrs[i]
- - (uintptr_t) hostaddrs[i + 1];
- else
- cur_node.tgt_offset = tgt->tgt_start
- + tgt->list[i].offset;
- }
- else
- cur_node.tgt_offset = tgt->list[i].key->tgt->tgt_start
- + tgt->list[i].key->tgt_offset
- + tgt->list[i].offset;
+ cur_node.tgt_offset = gomp_map_val (tgt, hostaddrs, i);
/* FIXME: see above FIXME comment. */
devicep->host2dev_func (devicep->target_id,
(void *) (tgt->tgt_start
@@ -1213,7 +1260,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)
@@ -1231,7 +1278,45 @@ gomp_target_fallback (void (*fn) (void *
*thr = old_thr;
}
-/* Helper function of GOMP_target{,_41} routines. */
+/* Host fallback with firstprivate map-type handling. */
+
+static void
+gomp_target_fallback_firstprivate (void (*fn) (void *), size_t mapnum,
+ void **hostaddrs, size_t *sizes,
+ unsigned short *kinds)
+{
+ size_t i, tgt_align = 0, tgt_size = 0;
+ char *tgt = NULL;
+ for (i = 0; i < mapnum; i++)
+ if ((kinds[i] & 0xff) == GOMP_MAP_FIRSTPRIVATE)
+ {
+ size_t align = (size_t) 1 << (kinds[i] >> 8);
+ if (tgt_align < align)
+ tgt_align = align;
+ tgt_size = (tgt_size + align - 1) & ~(align - 1);
+ tgt_size += sizes[i];
+ }
+ if (tgt_align)
+ {
+ tgt = gomp_alloca (tgt_size + tgt_align - 1);
+ uintptr_t al = (uintptr_t) tgt & (tgt_align - 1);
+ if (al)
+ tgt += tgt_align - al;
+ tgt_size = 0;
+ for (i = 0; i < mapnum; i++)
+ if ((kinds[i] & 0xff) == GOMP_MAP_FIRSTPRIVATE)
+ {
+ size_t align = (size_t) 1 << (kinds[i] >> 8);
+ tgt_size = (tgt_size + align - 1) & ~(align - 1);
+ memcpy (tgt + tgt_size, hostaddrs[i], sizes[i]);
+ hostaddrs[i] = tgt + tgt_size;
+ tgt_size = tgt_size + sizes[i];
+ }
+ }
+ gomp_target_fallback (fn, hostaddrs);
+}
+
+/* Helper function of GOMP_target{,_ext} routines. */
static void *
gomp_get_target_fn_addr (struct gomp_device_descr *devicep,
@@ -1294,13 +1379,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
@@ -1315,35 +1418,7 @@ GOMP_target_41 (int device, void (*fn) (
if (devicep == NULL
|| !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400))
{
- size_t i, tgt_align = 0, tgt_size = 0;
- char *tgt = NULL;
- for (i = 0; i < mapnum; i++)
- if ((kinds[i] & 0xff) == GOMP_MAP_FIRSTPRIVATE)
- {
- size_t align = (size_t) 1 << (kinds[i] >> 8);
- if (tgt_align < align)
- tgt_align = align;
- tgt_size = (tgt_size + align - 1) & ~(align - 1);
- tgt_size += sizes[i];
- }
- if (tgt_align)
- {
- tgt = gomp_alloca (tgt_size + tgt_align - 1);
- uintptr_t al = (uintptr_t) tgt & (tgt_align - 1);
- if (al)
- tgt += tgt_align - al;
- tgt_size = 0;
- for (i = 0; i < mapnum; i++)
- if ((kinds[i] & 0xff) == GOMP_MAP_FIRSTPRIVATE)
- {
- size_t align = (size_t) 1 << (kinds[i] >> 8);
- tgt_size = (tgt_size + align - 1) & ~(align - 1);
- memcpy (tgt + tgt_size, hostaddrs[i], sizes[i]);
- hostaddrs[i] = tgt + tgt_size;
- tgt_size = tgt_size + sizes[i];
- }
- }
- gomp_target_fallback (fn, hostaddrs);
+ gomp_target_fallback_firstprivate (fn, mapnum, hostaddrs, sizes, kinds);
return;
}
@@ -1366,7 +1441,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)
@@ -1405,8 +1480,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);
@@ -1448,9 +1523,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);
@@ -1527,7 +1602,7 @@ gomp_exit_data (struct gomp_device_descr
cur_node.host_end = cur_node.host_start + sizes[i];
splay_tree_key k = (kind == GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION
|| kind == GOMP_MAP_ZERO_LEN_ARRAY_SECTION)
- ? gomp_map_lookup (&devicep->mem_map, &cur_node)
+ ? gomp_map_0len_lookup (&devicep->mem_map, &cur_node)
: splay_tree_lookup (&devicep->mem_map, &cur_node);
if (!k)
continue;
@@ -1642,7 +1717,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))
@@ -1759,7 +1834,7 @@ omp_target_is_present (void *ptr, int de
cur_node.host_start = (uintptr_t) ptr;
cur_node.host_end = cur_node.host_start;
- splay_tree_key n = gomp_map_lookup (mem_map, &cur_node);
+ splay_tree_key n = gomp_map_0len_lookup (mem_map, &cur_node);
int ret = n != NULL;
gomp_mutex_unlock (&devicep->lock);
return ret;
@@ -41,7 +41,7 @@ main ()
if (omp_target_is_present (q, d) != 1
|| omp_target_is_present (&q[32], d) != 1
- || omp_target_is_present (&q[128], d) != 1)
+ || omp_target_is_present (&q[127], d) != 1)
abort ();
if (omp_target_memcpy (p, q, 128 * sizeof (int), sizeof (int), 0,
@@ -37,58 +37,6 @@ foo (int n)
}
if (err)
abort ();
- int on = n;
- #pragma omp target firstprivate (n) map(tofrom: n)
- {
- n++;
- }
- if (on != n)
- abort ();
- #pragma omp target map(tofrom: n) private (n)
- {
- n = 25;
- }
- if (on != n)
- abort ();
- for (i = 0; i < n; i++)
- a[i] += i;
- #pragma omp target map(to:a) firstprivate (a) map(from:err) private(i)
- {
- err = 0;
- for (i = 0; i < n; i++)
- if (a[i] != 8 * i)
- err = 1;
- }
- if (err)
- abort ();
- for (i = 0; i < n; i++)
- a[i] += i;
- #pragma omp target firstprivate (a) map(to:a) map(from:err) private(i)
- {
- err = 0;
- for (i = 0; i < n; i++)
- if (a[i] != 9 * i)
- err = 1;
- }
- if (err)
- abort ();
- for (i = 0; i < n; i++)
- a[i] += i;
- #pragma omp target map(tofrom:a) map(from:err) private(a, i)
- {
- err = 0;
- for (i = 0; i < n; i++)
- a[i] = 7;
- #pragma omp parallel for reduction(|:err)
- for (i = 0; i < n; i++)
- if (a[i] != 7)
- err |= 1;
- }
- if (err)
- abort ();
- for (i = 0; i < n; i++)
- if (a[i] != 10 * i)
- abort ();
}
int
@@ -1,21 +1,29 @@
extern void abort (void);
-void
+__attribute__((noinline, noclone)) void
foo (int *p, int *q, int *r, int n, int m)
{
int i, err, *s = r;
+ int sep = 1;
+ #pragma omp target map(to:sep)
+ sep = 0;
#pragma omp target data map(to:p[0:8])
{
/* For zero length array sections, p points to the start of
- already mapped range, q to the end of it, and r does not point
- to an mapped range. */
+ already mapped range, q to the end of it (with nothing mapped
+ after it), and r does not point to an mapped range. */
#pragma omp target map(alloc:p[:0]) map(to:q[:0]) map(from:r[:0]) private(i) map(from:err) firstprivate (s)
{
err = 0;
for (i = 0; i < 8; i++)
- if (p[i] != i + 1 || q[i - 8] != i + 1)
+ if (p[i] != i + 1)
err = 1;
- if (p + 8 != q || (r != (int *) 0 && r != s))
+ if (sep)
+ {
+ if (q != (int *) 0 || r != (int *) 0)
+ err = 1;
+ }
+ else if (p + 8 != q || r != s)
err = 1;
}
if (err)
@@ -25,9 +33,14 @@ foo (int *p, int *q, int *r, int n, int
{
err = 0;
for (i = 0; i < 8; i++)
- if (p[i] != i + 1 || q[i - 8] != i + 1)
+ if (p[i] != i + 1)
err = 1;
- if (p + 8 != q || (r != (int *) 0 && r != s))
+ if (sep)
+ {
+ if (q != (int *) 0 || r != (int *) 0)
+ err = 1;
+ }
+ else if (p + 8 != q || r != s)
err = 1;
}
if (err)
@@ -38,9 +51,14 @@ foo (int *p, int *q, int *r, int n, int
{
err = 0;
for (i = 0; i < 8; i++)
- if (p[i] != i + 1 || q[i - 8] != i + 1)
+ if (p[i] != i + 1)
err = 1;
- if (p + 8 != q || (r != (int *) 0 && r != s))
+ if (sep)
+ {
+ if (q != (int *) 0 || r != (int *) 0)
+ err = 1;
+ }
+ else if (p + 8 != q || r != s)
err = 1;
}
if (err)
@@ -69,7 +87,14 @@ foo (int *p, int *q, int *r, int n, int
for (i = 0; i < 8; i++)
if (p[i] != i + 1)
err = 1;
- if (q[0] != 9 || r != q + 1)
+ if (q[0] != 9)
+ err = 1;
+ else if (sep)
+ {
+ if (r != (int *) 0)
+ err = 1;
+ }
+ else if (r != q + 1)
err = 1;
}
if (err)
@@ -81,7 +106,14 @@ foo (int *p, int *q, int *r, int n, int
for (i = 0; i < 8; i++)
if (p[i] != i + 1)
err = 1;
- if (q[0] != 9 || r != q + 1)
+ if (q[0] != 9)
+ err = 1;
+ else if (sep)
+ {
+ if (r != (int *) 0)
+ err = 1;
+ }
+ else if (r != q + 1)
err = 1;
}
if (err)
@@ -94,7 +126,14 @@ foo (int *p, int *q, int *r, int n, int
for (i = 0; i < 8; i++)
if (p[i] != i + 1)
err = 1;
- if (q[0] != 9 || r != q + 1)
+ if (q[0] != 9)
+ err = 1;
+ else if (sep)
+ {
+ if (r != (int *) 0)
+ err = 1;
+ }
+ else if (r != q + 1)
err = 1;
}
if (err)