X-Git-Url: https://git.libre-soc.org/?a=blobdiff_plain;f=gcc%2Fomp-low.c;h=09a8cbdc433ab9ec6be9dc2d728470cf941c9ba0;hb=65eee57a8cccc77a1bfd5ad5cde53460ad564124;hp=a855c5b2f8b959e861ff5137ac0ff7d8c8ccc4b6;hpb=6c1dae73cd2ceb6a326f786f6d4e8674863de45e;p=gcc.git diff --git a/gcc/omp-low.c b/gcc/omp-low.c index a855c5b2f8b..09a8cbdc433 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -4,7 +4,7 @@ Contributed by Diego Novillo - Copyright (C) 2005-2019 Free Software Foundation, Inc. + Copyright (C) 2005-2020 Free Software Foundation, Inc. This file is part of GCC. @@ -50,14 +50,13 @@ along with GCC; see the file COPYING3. If not see #include "splay-tree.h" #include "omp-general.h" #include "omp-low.h" -#include "omp-grid.h" #include "gimple-low.h" +#include "alloc-pool.h" #include "symbol-summary.h" #include "tree-nested.h" #include "context.h" #include "gomp-constants.h" #include "gimple-pretty-print.h" -#include "hsa-common.h" #include "stringpool.h" #include "attribs.h" @@ -127,6 +126,22 @@ struct omp_context corresponding tracking loop iteration variables. */ hash_map *lastprivate_conditional_map; + /* And a hash map from the allocate variables to their corresponding + allocators. */ + hash_map *allocate_map; + + /* A tree_list of the reduction clauses in this context. This is + only used for checking the consistency of OpenACC reduction + clauses in scan_omp_for and is not guaranteed to contain a valid + value outside of this function. */ + tree local_reduction_clauses; + + /* A tree_list of the reduction clauses in outer contexts. This is + only used for checking the consistency of OpenACC reduction + clauses in scan_omp_for and is not guaranteed to contain a valid + value outside of this function. */ + tree outer_reduction_clauses; + /* Nesting depth of this context. Used to beautify error messages re invalid gotos. The outermost ctx is depth 1, with depth 0 being reserved for the main body of the function. */ @@ -150,12 +165,19 @@ struct omp_context /* True in the second simd loop of for simd with inscan reductions. */ bool for_simd_scan_phase; + + /* True if there is order(concurrent) clause on the construct. */ + bool order_concurrent; + + /* True if there is bind clause on the construct (i.e. a loop construct). */ + bool loop_p; }; static splay_tree all_contexts; static int taskreg_nesting_level; static int target_nesting_level; static bitmap task_shared_vars; +static bitmap global_nonaddressable_vars; static vec taskreg_contexts; static void scan_omp (gimple_seq *, omp_context *); @@ -171,18 +193,22 @@ static tree scan_omp_1_op (tree *, int *, void *); *handled_ops_p = false; \ break; -/* Return true if CTX corresponds to an oacc parallel region. */ +/* Return whether CTX represents an OpenACC 'parallel' or 'serial' construct. + (This doesn't include OpenACC 'kernels' decomposed parts.) */ static bool -is_oacc_parallel (omp_context *ctx) +is_oacc_parallel_or_serial (omp_context *ctx) { enum gimple_code outer_type = gimple_code (ctx->stmt); return ((outer_type == GIMPLE_OMP_TARGET) - && (gimple_omp_target_kind (ctx->stmt) - == GF_OMP_TARGET_KIND_OACC_PARALLEL)); + && ((gimple_omp_target_kind (ctx->stmt) + == GF_OMP_TARGET_KIND_OACC_PARALLEL) + || (gimple_omp_target_kind (ctx->stmt) + == GF_OMP_TARGET_KIND_OACC_SERIAL))); } -/* Return true if CTX corresponds to an oacc kernels region. */ +/* Return whether CTX represents an OpenACC 'kernels' construct. + (This doesn't include OpenACC 'kernels' decomposed parts.) */ static bool is_oacc_kernels (omp_context *ctx) @@ -193,6 +219,36 @@ is_oacc_kernels (omp_context *ctx) == GF_OMP_TARGET_KIND_OACC_KERNELS)); } +/* Return whether CTX represents an OpenACC 'kernels' decomposed part. */ + +static bool +is_oacc_kernels_decomposed_part (omp_context *ctx) +{ + enum gimple_code outer_type = gimple_code (ctx->stmt); + return ((outer_type == GIMPLE_OMP_TARGET) + && ((gimple_omp_target_kind (ctx->stmt) + == GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_PARALLELIZED) + || (gimple_omp_target_kind (ctx->stmt) + == GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_GANG_SINGLE) + || (gimple_omp_target_kind (ctx->stmt) + == GF_OMP_TARGET_KIND_OACC_DATA_KERNELS))); +} + +/* Return true if STMT corresponds to an OpenMP target region. */ +static bool +is_omp_target (gimple *stmt) +{ + if (gimple_code (stmt) == GIMPLE_OMP_TARGET) + { + int kind = gimple_omp_target_kind (stmt); + return (kind == GF_OMP_TARGET_KIND_REGION + || kind == GF_OMP_TARGET_KIND_DATA + || kind == GF_OMP_TARGET_KIND_ENTER_DATA + || kind == GF_OMP_TARGET_KIND_EXIT_DATA); + } + return false; +} + /* If DECL is the artificial dummy VAR_DECL created for non-static data member privatization, return the underlying "this" parameter, otherwise return NULL. */ @@ -420,7 +476,26 @@ use_pointer_for_field (tree decl, omp_context *shared_ctx) /* Do not use copy-in/copy-out for variables that have their address taken. */ - if (TREE_ADDRESSABLE (decl)) + if (is_global_var (decl)) + { + /* For file scope vars, track whether we've seen them as + non-addressable initially and in that case, keep the same + answer for the duration of the pass, even when they are made + addressable later on e.g. through reduction expansion. Global + variables which weren't addressable before the pass will not + have their privatized copies address taken. See PR91216. */ + if (!TREE_ADDRESSABLE (decl)) + { + if (!global_nonaddressable_vars) + global_nonaddressable_vars = BITMAP_ALLOC (NULL); + bitmap_set_bit (global_nonaddressable_vars, DECL_UID (decl)); + } + else if (!global_nonaddressable_vars + || !bitmap_bit_p (global_nonaddressable_vars, + DECL_UID (decl))) + return true; + } + else if (TREE_ADDRESSABLE (decl)) return true; /* lower_send_shared_vars only uses copy-in, but not copy-out @@ -441,18 +516,30 @@ use_pointer_for_field (tree decl, omp_context *shared_ctx) omp_context *up; for (up = shared_ctx->outer; up; up = up->outer) - if (is_taskreg_ctx (up) && maybe_lookup_decl (decl, up)) + if ((is_taskreg_ctx (up) + || (gimple_code (up->stmt) == GIMPLE_OMP_TARGET + && is_gimple_omp_offloaded (up->stmt))) + && maybe_lookup_decl (decl, up)) break; if (up) { tree c; - for (c = gimple_omp_taskreg_clauses (up->stmt); - c; c = OMP_CLAUSE_CHAIN (c)) - if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_SHARED - && OMP_CLAUSE_DECL (c) == decl) - break; + if (gimple_code (up->stmt) == GIMPLE_OMP_TARGET) + { + for (c = gimple_omp_target_clauses (up->stmt); + c; c = OMP_CLAUSE_CHAIN (c)) + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_DECL (c) == decl) + break; + } + else + for (c = gimple_omp_taskreg_clauses (up->stmt); + c; c = OMP_CLAUSE_CHAIN (c)) + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_SHARED + && OMP_CLAUSE_DECL (c) == decl) + break; if (c) goto maybe_mark_addressable_and_ret; @@ -498,8 +585,10 @@ omp_copy_decl_2 (tree var, tree name, tree type, omp_context *ctx) it's address. But we don't need to take address of privatizations from that var. */ if (TREE_ADDRESSABLE (var) - && task_shared_vars - && bitmap_bit_p (task_shared_vars, DECL_UID (var))) + && ((task_shared_vars + && bitmap_bit_p (task_shared_vars, DECL_UID (var))) + || (global_nonaddressable_vars + && bitmap_bit_p (global_nonaddressable_vars, DECL_UID (var))))) TREE_ADDRESSABLE (copy) = 0; ctx->block_vars = copy; @@ -577,7 +666,8 @@ build_outer_var_ref (tree var, omp_context *ctx, x = build_receiver_ref (var, by_ref, ctx); } else if ((gimple_code (ctx->stmt) == GIMPLE_OMP_FOR - && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_SIMD) + && gimple_omp_for_kind (ctx->stmt) == GF_OMP_FOR_KIND_SIMD) + || ctx->loop_p || (code == OMP_CLAUSE_PRIVATE && (gimple_code (ctx->stmt) == GIMPLE_OMP_FOR || gimple_code (ctx->stmt) == GIMPLE_OMP_SECTIONS @@ -624,15 +714,7 @@ build_outer_var_ref (tree var, omp_context *ctx, } } else if (outer) - { - if (gimple_code (outer->stmt) == GIMPLE_OMP_GRID_BODY) - { - outer = outer->outer; - gcc_assert (outer - && gimple_code (outer->stmt) != GIMPLE_OMP_GRID_BODY); - } - x = lookup_decl (var, outer); - } + x = lookup_decl (var, outer); else if (omp_is_reference (var)) /* This can happen with orphaned constructs. If var is reference, it is possible it is shared and as such valid. */ @@ -686,6 +768,11 @@ install_var_field (tree var, bool by_ref, int mask, omp_context *ctx) tree field, type, sfield = NULL_TREE; splay_tree_key key = (splay_tree_key) var; + if ((mask & 16) != 0) + { + key = (splay_tree_key) &DECL_NAME (var); + gcc_checking_assert (key != (splay_tree_key) var); + } if ((mask & 8) != 0) { key = (splay_tree_key) &DECL_UID (var); @@ -699,6 +786,9 @@ install_var_field (tree var, bool by_ref, int mask, omp_context *ctx) || !is_gimple_omp_oacc (ctx->stmt)); type = TREE_TYPE (var); + if ((mask & 16) != 0) + type = lang_hooks.decls.omp_array_data (var, true); + /* Prevent redeclaring the var in the split-off function with a restrict pointer type. Note that we only clear type itself, restrict qualifiers in the pointed-to type will be ignored by points-to analysis. */ @@ -713,7 +803,7 @@ install_var_field (tree var, bool by_ref, int mask, omp_context *ctx) } else if (by_ref) type = build_pointer_type (type); - else if ((mask & 3) == 1 && omp_is_reference (var)) + else if ((mask & (32 | 3)) == 1 && omp_is_reference (var)) type = TREE_TYPE (type); field = build_decl (DECL_SOURCE_LOCATION (var), @@ -723,7 +813,7 @@ install_var_field (tree var, bool by_ref, int mask, omp_context *ctx) side effect of making dwarf2out ignore this member, so for helpful debugging we clear it later in delete_omp_context. */ DECL_ABSTRACT_ORIGIN (field) = var; - if (type == TREE_TYPE (var)) + if ((mask & 16) == 0 && type == TREE_TYPE (var)) { SET_DECL_ALIGN (field, DECL_ALIGN (var)); DECL_USER_ALIGN (field) = DECL_USER_ALIGN (var); @@ -973,6 +1063,7 @@ delete_omp_context (splay_tree_value value) } delete ctx->lastprivate_conditional_map; + delete ctx->allocate_map; XDELETE (ctx); } @@ -1044,6 +1135,20 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) tree c, decl; bool scan_array_reductions = false; + for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_ALLOCATE + && (OMP_CLAUSE_ALLOCATE_ALLOCATOR (c) == NULL_TREE + /* omp_default_mem_alloc is 1 */ + || !integer_onep (OMP_CLAUSE_ALLOCATE_ALLOCATOR (c)))) + { + if (ctx->allocate_map == NULL) + ctx->allocate_map = new hash_map; + ctx->allocate_map->put (OMP_CLAUSE_DECL (c), + OMP_CLAUSE_ALLOCATE_ALLOCATOR (c) + ? OMP_CLAUSE_ALLOCATE_ALLOCATOR (c) + : integer_zero_node); + } + for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) { bool by_ref; @@ -1060,6 +1165,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_SHARED: decl = OMP_CLAUSE_DECL (c); + if (ctx->allocate_map && ctx->allocate_map->get (decl)) + ctx->allocate_map->remove (decl); /* Ignore shared directives in teams construct inside of target construct. */ if (gimple_code (ctx->stmt) == GIMPLE_OMP_TEAMS @@ -1101,8 +1208,32 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) goto do_private; case OMP_CLAUSE_REDUCTION: + /* Collect 'reduction' clauses on OpenACC compute construct. */ + if (is_gimple_omp_oacc (ctx->stmt) + && is_gimple_omp_offloaded (ctx->stmt)) + { + /* No 'reduction' clauses on OpenACC 'kernels'. */ + gcc_checking_assert (!is_oacc_kernels (ctx)); + /* Likewise, on OpenACC 'kernels' decomposed parts. */ + gcc_checking_assert (!is_oacc_kernels_decomposed_part (ctx)); + + ctx->local_reduction_clauses + = tree_cons (NULL, c, ctx->local_reduction_clauses); + } + /* FALLTHRU */ + case OMP_CLAUSE_IN_REDUCTION: decl = OMP_CLAUSE_DECL (c); + if (ctx->allocate_map + && ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION + && (OMP_CLAUSE_REDUCTION_INSCAN (c) + || OMP_CLAUSE_REDUCTION_TASK (c))) + || is_task_ctx (ctx))) + { + /* For now. */ + if (ctx->allocate_map->get (decl)) + ctx->allocate_map->remove (decl); + } if (TREE_CODE (decl) == MEM_REF) { tree t = TREE_OPERAND (decl, 0); @@ -1186,7 +1317,16 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) if (is_variable_sized (decl)) { if (is_task_ctx (ctx)) - install_var_field (decl, false, 1, ctx); + { + if (ctx->allocate_map + && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE) + { + /* For now. */ + if (ctx->allocate_map->get (decl)) + ctx->allocate_map->remove (decl); + } + install_var_field (decl, false, 1, ctx); + } break; } else if (is_taskreg_ctx (ctx)) @@ -1198,7 +1338,11 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) if (is_task_ctx (ctx) && (global || by_ref || omp_is_reference (decl))) { - install_var_field (decl, false, 1, ctx); + if (ctx->allocate_map + && ctx->allocate_map->get (decl)) + install_var_field (decl, by_ref, 32 | 1, ctx); + else + install_var_field (decl, false, 1, ctx); if (!global) install_var_field (decl, by_ref, 2, ctx); } @@ -1209,11 +1353,19 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) break; case OMP_CLAUSE_USE_DEVICE_PTR: + case OMP_CLAUSE_USE_DEVICE_ADDR: decl = OMP_CLAUSE_DECL (c); - if (TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) - install_var_field (decl, true, 3, ctx); + + /* Fortran array descriptors. */ + if (lang_hooks.decls.omp_array_data (decl, true)) + install_var_field (decl, false, 19, ctx); + else if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR + && !omp_is_reference (decl) + && !omp_is_allocatable_or_ptr (decl)) + || TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) + install_var_field (decl, true, 11, ctx); else - install_var_field (decl, false, 3, ctx); + install_var_field (decl, false, 11, ctx); if (DECL_SIZE (decl) && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST) { @@ -1278,11 +1430,14 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) && DECL_P (decl) && ((OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER && (OMP_CLAUSE_MAP_KIND (c) - != GOMP_MAP_FIRSTPRIVATE_REFERENCE)) + != GOMP_MAP_FIRSTPRIVATE_REFERENCE) + && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH + && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH) || TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_TO && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_FROM && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_TOFROM + && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET && is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx)) && varpool_node::get_create (decl)->offloadable && !lookup_attribute ("omp declare target link", @@ -1298,6 +1453,40 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)) break; } + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && DECL_P (decl) + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH) + && is_omp_target (ctx->stmt)) + { + /* If this is an offloaded region, an attach operation should + only exist when the pointer variable is mapped in a prior + clause. */ + if (is_gimple_omp_offloaded (ctx->stmt)) + gcc_assert + (maybe_lookup_decl (decl, ctx) + || (is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx)) + && lookup_attribute ("omp declare target", + DECL_ATTRIBUTES (decl)))); + + /* By itself, attach/detach is generated as part of pointer + variable mapping and should not create new variables in the + offloaded region, however sender refs for it must be created + for its address to be passed to the runtime. */ + tree field + = build_decl (OMP_CLAUSE_LOCATION (c), + FIELD_DECL, NULL_TREE, ptr_type_node); + SET_DECL_ALIGN (field, TYPE_ALIGN (ptr_type_node)); + insert_field_into_struct (ctx->record_type, field); + /* To not clash with a map of the pointer variable itself, + attach/detach maps have their field looked up by the *clause* + tree expression, not the decl. */ + gcc_assert (!splay_tree_lookup (ctx->field_map, + (splay_tree_key) c)); + splay_tree_insert (ctx->field_map, (splay_tree_key) c, + (splay_tree_value) field); + break; + } if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER || (OMP_CLAUSE_MAP_KIND (c) @@ -1382,12 +1571,12 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) } break; - case OMP_CLAUSE__GRIDDIM_: - if (ctx->outer) - { - scan_omp_op (&OMP_CLAUSE__GRIDDIM__SIZE (c), ctx->outer); - scan_omp_op (&OMP_CLAUSE__GRIDDIM__GROUP (c), ctx->outer); - } + case OMP_CLAUSE_ORDER: + ctx->order_concurrent = true; + break; + + case OMP_CLAUSE_BIND: + ctx->loop_p = true; break; case OMP_CLAUSE_NOWAIT: @@ -1417,6 +1606,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_IF_PRESENT: case OMP_CLAUSE_FINALIZE: case OMP_CLAUSE_TASK_REDUCTION: + case OMP_CLAUSE_ALLOCATE: break; case OMP_CLAUSE_ALIGNED: @@ -1434,7 +1624,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) install_var_local (decl, ctx); } else if (gimple_code (ctx->stmt) == GIMPLE_OMP_FOR - && (gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_SIMD) + && gimple_omp_for_kind (ctx->stmt) == GF_OMP_FOR_KIND_SIMD && !OMP_CLAUSE__CONDTEMP__ITER (c)) install_var_local (decl, ctx); break; @@ -1537,6 +1727,11 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) && is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx)) && varpool_node::get_create (decl)->offloadable) break; + if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH) + && is_omp_target (ctx->stmt) + && !is_gimple_omp_offloaded (ctx->stmt)) + break; if (DECL_P (decl)) { if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER @@ -1584,6 +1779,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_SIMDLEN: case OMP_CLAUSE_ALIGNED: case OMP_CLAUSE_DEPEND: + case OMP_CLAUSE_ALLOCATE: case OMP_CLAUSE__LOOPTEMP_: case OMP_CLAUSE__REDUCTEMP_: case OMP_CLAUSE_TO: @@ -1595,7 +1791,10 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_SIMD: case OMP_CLAUSE_NOGROUP: case OMP_CLAUSE_DEFAULTMAP: + case OMP_CLAUSE_ORDER: + case OMP_CLAUSE_BIND: case OMP_CLAUSE_USE_DEVICE_PTR: + case OMP_CLAUSE_USE_DEVICE_ADDR: case OMP_CLAUSE_NONTEMPORAL: case OMP_CLAUSE_ASYNC: case OMP_CLAUSE_WAIT: @@ -1609,7 +1808,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_AUTO: case OMP_CLAUSE_SEQ: case OMP_CLAUSE_TILE: - case OMP_CLAUSE__GRIDDIM_: case OMP_CLAUSE__SIMT_: case OMP_CLAUSE_IF_PRESENT: case OMP_CLAUSE_FINALIZE: @@ -1848,12 +2046,38 @@ add_taskreg_looptemp_clauses (enum gf_mask msk, gimple *stmt, GIMPLE_OMP_FOR, add one more temporaries for the total number of iterations (product of count1 ... countN-1). */ if (omp_find_clause (gimple_omp_for_clauses (for_stmt), - OMP_CLAUSE_LASTPRIVATE)) - count++; - else if (msk == GF_OMP_FOR_KIND_FOR - && omp_find_clause (gimple_omp_parallel_clauses (stmt), - OMP_CLAUSE_LASTPRIVATE)) - count++; + OMP_CLAUSE_LASTPRIVATE) + || (msk == GF_OMP_FOR_KIND_FOR + && omp_find_clause (gimple_omp_parallel_clauses (stmt), + OMP_CLAUSE_LASTPRIVATE))) + { + tree temp = create_tmp_var (type); + tree c = build_omp_clause (UNKNOWN_LOCATION, + OMP_CLAUSE__LOOPTEMP_); + insert_decl_map (&outer_ctx->cb, temp, temp); + OMP_CLAUSE_DECL (c) = temp; + OMP_CLAUSE_CHAIN (c) = gimple_omp_taskreg_clauses (stmt); + gimple_omp_taskreg_set_clauses (stmt, c); + } + if (fd.non_rect + && fd.last_nonrect == fd.first_nonrect + 1) + if (tree v = gimple_omp_for_index (for_stmt, fd.last_nonrect)) + if (!TYPE_UNSIGNED (TREE_TYPE (v))) + { + v = gimple_omp_for_index (for_stmt, fd.first_nonrect); + tree type2 = TREE_TYPE (v); + count++; + for (i = 0; i < 3; i++) + { + tree temp = create_tmp_var (type2); + tree c = build_omp_clause (UNKNOWN_LOCATION, + OMP_CLAUSE__LOOPTEMP_); + insert_decl_map (&outer_ctx->cb, temp, temp); + OMP_CLAUSE_DECL (c) = temp; + OMP_CLAUSE_CHAIN (c) = gimple_omp_taskreg_clauses (stmt); + gimple_omp_taskreg_set_clauses (stmt, c); + } + } } for (i = 0; i < count; i++) { @@ -1932,11 +2156,8 @@ scan_omp_parallel (gimple_stmt_iterator *gsi, omp_context *outer_ctx) DECL_NAMELESS (name) = 1; TYPE_NAME (ctx->record_type) = name; TYPE_ARTIFICIAL (ctx->record_type) = 1; - if (!gimple_omp_parallel_grid_phony (stmt)) - { - create_omp_child_function (ctx, false); - gimple_omp_parallel_set_child_fn (stmt, ctx->cb.dst_fn); - } + create_omp_child_function (ctx, false); + gimple_omp_parallel_set_child_fn (stmt, ctx->cb.dst_fn); scan_sharing_clauses (gimple_omp_parallel_clauses (stmt), ctx); scan_omp (gimple_omp_body_ptr (stmt), ctx); @@ -2225,7 +2446,9 @@ enclosing_target_ctx (omp_context *ctx) return ctx; } -/* Return true if ctx is part of an oacc kernels region. */ +/* Return whether CTX's parent compute construct is an OpenACC 'kernels' + construct. + (This doesn't include OpenACC 'kernels' decomposed parts.) */ static bool ctx_in_oacc_kernels_region (omp_context *ctx) @@ -2241,7 +2464,8 @@ ctx_in_oacc_kernels_region (omp_context *ctx) return false; } -/* Check the parallelism clauses inside a kernels regions. +/* Check the parallelism clauses inside a OpenACC 'kernels' region. + (This doesn't include OpenACC 'kernels' decomposed parts.) Until kernels handling moves to use the same loop indirection scheme as parallel, we need to do this checking early. */ @@ -2319,35 +2543,133 @@ scan_omp_for (gomp_for *stmt, omp_context *outer_ctx) { omp_context *tgt = enclosing_target_ctx (outer_ctx); - if (!tgt || is_oacc_parallel (tgt)) + if (!(tgt && is_oacc_kernels (tgt))) for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) { - char const *check = NULL; - + tree c_op0; switch (OMP_CLAUSE_CODE (c)) { case OMP_CLAUSE_GANG: - check = "gang"; + c_op0 = OMP_CLAUSE_GANG_EXPR (c); break; case OMP_CLAUSE_WORKER: - check = "worker"; + c_op0 = OMP_CLAUSE_WORKER_EXPR (c); break; case OMP_CLAUSE_VECTOR: - check = "vector"; + c_op0 = OMP_CLAUSE_VECTOR_EXPR (c); break; default: - break; + continue; } - if (check && OMP_CLAUSE_OPERAND (c, 0)) - error_at (gimple_location (stmt), - "argument not permitted on %qs clause in" - " OpenACC %", check); + if (c_op0) + { + /* By construction, this is impossible for OpenACC 'kernels' + decomposed parts. */ + gcc_assert (!(tgt && is_oacc_kernels_decomposed_part (tgt))); + + error_at (OMP_CLAUSE_LOCATION (c), + "argument not permitted on %qs clause", + omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + if (tgt) + inform (gimple_location (tgt->stmt), + "enclosing parent compute construct"); + else if (oacc_get_fn_attrib (current_function_decl)) + inform (DECL_SOURCE_LOCATION (current_function_decl), + "enclosing routine"); + else + gcc_unreachable (); + } } + if (tgt && is_oacc_kernels (tgt)) + check_oacc_kernel_gwv (stmt, ctx); + + /* Collect all variables named in reductions on this loop. Ensure + that, if this loop has a reduction on some variable v, and there is + a reduction on v somewhere in an outer context, then there is a + reduction on v on all intervening loops as well. */ + tree local_reduction_clauses = NULL; + for (tree c = gimple_omp_for_clauses (stmt); c; c = OMP_CLAUSE_CHAIN (c)) + { + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION) + local_reduction_clauses + = tree_cons (NULL, c, local_reduction_clauses); + } + if (ctx->outer_reduction_clauses == NULL && ctx->outer != NULL) + ctx->outer_reduction_clauses + = chainon (unshare_expr (ctx->outer->local_reduction_clauses), + ctx->outer->outer_reduction_clauses); + tree outer_reduction_clauses = ctx->outer_reduction_clauses; + tree local_iter = local_reduction_clauses; + for (; local_iter; local_iter = TREE_CHAIN (local_iter)) + { + tree local_clause = TREE_VALUE (local_iter); + tree local_var = OMP_CLAUSE_DECL (local_clause); + tree_code local_op = OMP_CLAUSE_REDUCTION_CODE (local_clause); + bool have_outer_reduction = false; + tree ctx_iter = outer_reduction_clauses; + for (; ctx_iter; ctx_iter = TREE_CHAIN (ctx_iter)) + { + tree outer_clause = TREE_VALUE (ctx_iter); + tree outer_var = OMP_CLAUSE_DECL (outer_clause); + tree_code outer_op = OMP_CLAUSE_REDUCTION_CODE (outer_clause); + if (outer_var == local_var && outer_op != local_op) + { + warning_at (OMP_CLAUSE_LOCATION (local_clause), 0, + "conflicting reduction operations for %qE", + local_var); + inform (OMP_CLAUSE_LOCATION (outer_clause), + "location of the previous reduction for %qE", + outer_var); + } + if (outer_var == local_var) + { + have_outer_reduction = true; + break; + } + } + if (have_outer_reduction) + { + /* There is a reduction on outer_var both on this loop and on + some enclosing loop. Walk up the context tree until such a + loop with a reduction on outer_var is found, and complain + about all intervening loops that do not have such a + reduction. */ + struct omp_context *curr_loop = ctx->outer; + bool found = false; + while (curr_loop != NULL) + { + tree curr_iter = curr_loop->local_reduction_clauses; + for (; curr_iter; curr_iter = TREE_CHAIN (curr_iter)) + { + tree curr_clause = TREE_VALUE (curr_iter); + tree curr_var = OMP_CLAUSE_DECL (curr_clause); + if (curr_var == local_var) + { + found = true; + break; + } + } + if (!found) + warning_at (gimple_location (curr_loop->stmt), 0, + "nested loop in reduction needs " + "reduction clause for %qE", + local_var); + else + break; + curr_loop = curr_loop->outer; + } + } + } + ctx->local_reduction_clauses = local_reduction_clauses; + ctx->outer_reduction_clauses + = chainon (unshare_expr (ctx->local_reduction_clauses), + ctx->outer_reduction_clauses); + if (tgt && is_oacc_kernels (tgt)) { /* Strip out reductions, as they are not handled yet. */ @@ -2356,7 +2678,7 @@ scan_omp_for (gomp_for *stmt, omp_context *outer_ctx) while (tree probe = *prev_ptr) { tree *next_ptr = &OMP_CLAUSE_CHAIN (probe); - + if (OMP_CLAUSE_CODE (probe) == OMP_CLAUSE_REDUCTION) *prev_ptr = *next_ptr; else @@ -2364,7 +2686,6 @@ scan_omp_for (gomp_for *stmt, omp_context *outer_ctx) } gimple_omp_for_set_clauses (stmt, clauses); - check_oacc_kernel_gwv (stmt, ctx); } } @@ -2630,11 +2951,6 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx) { tree c; - if (ctx && gimple_code (ctx->stmt) == GIMPLE_OMP_GRID_BODY) - /* GRID_BODY is an artificial construct, nesting rules will be checked in - the original copy of its contents. */ - return true; - /* No nesting of non-OpenACC STMT (that is, an OpenMP one, or a GOMP builtin) inside an OpenACC CTX. */ if (!(is_gimple_omp (stmt) @@ -2667,9 +2983,21 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx) && gimple_code (ctx->outer->stmt) == GIMPLE_OMP_FOR) ctx = ctx->outer; if (gimple_code (ctx->stmt) == GIMPLE_OMP_FOR - && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_SIMD) + && gimple_omp_for_kind (ctx->stmt) == GF_OMP_FOR_KIND_SIMD + && !ctx->loop_p) { c = NULL_TREE; + if (ctx->order_concurrent + && (gimple_code (stmt) == GIMPLE_OMP_ORDERED + || gimple_code (stmt) == GIMPLE_OMP_ATOMIC_LOAD + || gimple_code (stmt) == GIMPLE_OMP_ATOMIC_STORE)) + { + error_at (gimple_location (stmt), + "OpenMP constructs other than %, %" + " or % may not be nested inside a region with" + " the % clause"); + return false; + } if (gimple_code (stmt) == GIMPLE_OMP_ORDERED) { c = gimple_omp_ordered_clauses (as_a (stmt)); @@ -2685,7 +3013,8 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx) { error_at (gimple_location (stmt), "% must be closely " - "nested inside of % region"); + "nested inside of %<%s simd%> region", + lang_GNU_Fortran () ? "do" : "for"); return false; } return true; @@ -2695,31 +3024,52 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx) || gimple_code (stmt) == GIMPLE_OMP_ATOMIC_STORE || gimple_code (stmt) == GIMPLE_OMP_SCAN) return true; + else if (gimple_code (stmt) == GIMPLE_OMP_FOR + && gimple_omp_for_kind (ctx->stmt) == GF_OMP_FOR_KIND_SIMD) + return true; error_at (gimple_location (stmt), - "OpenMP constructs other than %<#pragma omp ordered simd%>" - " or %<#pragma omp atomic%> may not be nested inside" - " % region"); + "OpenMP constructs other than " + "%, %, % or % may " + "not be nested inside % region"); return false; } else if (gimple_code (ctx->stmt) == GIMPLE_OMP_TEAMS) { if ((gimple_code (stmt) != GIMPLE_OMP_FOR - || ((gimple_omp_for_kind (stmt) != GF_OMP_FOR_KIND_DISTRIBUTE) - && (gimple_omp_for_kind (stmt) != GF_OMP_FOR_KIND_GRID_LOOP))) + || (gimple_omp_for_kind (stmt) != GF_OMP_FOR_KIND_DISTRIBUTE + && omp_find_clause (gimple_omp_for_clauses (stmt), + OMP_CLAUSE_BIND) == NULL_TREE)) && gimple_code (stmt) != GIMPLE_OMP_PARALLEL) { error_at (gimple_location (stmt), - "only % or % regions are " - "allowed to be strictly nested inside % " - "region"); + "only %, % or % " + "regions are allowed to be strictly nested inside " + "% region"); return false; } } + else if (ctx->order_concurrent + && gimple_code (stmt) != GIMPLE_OMP_PARALLEL + && (gimple_code (stmt) != GIMPLE_OMP_FOR + || gimple_omp_for_kind (stmt) != GF_OMP_FOR_KIND_SIMD) + && gimple_code (stmt) != GIMPLE_OMP_SCAN) + { + if (ctx->loop_p) + error_at (gimple_location (stmt), + "OpenMP constructs other than %, % or " + "% may not be nested inside a % region"); + else + error_at (gimple_location (stmt), + "OpenMP constructs other than %, % or " + "% may not be nested inside a region with " + "the % clause"); + return false; + } } switch (gimple_code (stmt)) { case GIMPLE_OMP_FOR: - if (gimple_omp_for_kind (stmt) & GF_OMP_FOR_SIMD) + if (gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_SIMD) return true; if (gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_DISTRIBUTE) { @@ -2735,6 +3085,11 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx) /* We split taskloop into task and nested taskloop in it. */ if (gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_TASKLOOP) return true; + /* For now, hope this will change and loop bind(parallel) will not + be allowed in lots of contexts. */ + if (gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_FOR + && omp_find_clause (gimple_omp_for_clauses (stmt), OMP_CLAUSE_BIND)) + return true; if (gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_OACC_LOOP) { bool ok = false; @@ -2752,6 +3107,9 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx) { case GF_OMP_TARGET_KIND_OACC_PARALLEL: case GF_OMP_TARGET_KIND_OACC_KERNELS: + case GF_OMP_TARGET_KIND_OACC_SERIAL: + case GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_PARALLELIZED: + case GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_GANG_SINGLE: ok = true; break; @@ -2785,8 +3143,8 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx) const char *construct = (DECL_FUNCTION_CODE (gimple_call_fndecl (stmt)) == BUILT_IN_GOMP_CANCEL) - ? "#pragma omp cancel" - : "#pragma omp cancellation point"; + ? "cancel" + : "cancellation point"; if (ctx == NULL) { error_at (gimple_location (stmt), "orphaned %qs construct", @@ -2799,7 +3157,7 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx) { case 1: if (gimple_code (ctx->stmt) != GIMPLE_OMP_PARALLEL) - bad = "#pragma omp parallel"; + bad = "parallel"; else if (DECL_FUNCTION_CODE (gimple_call_fndecl (stmt)) == BUILT_IN_GOMP_CANCEL && !integer_zerop (gimple_call_arg (stmt, 1))) @@ -2809,7 +3167,7 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx) case 2: if (gimple_code (ctx->stmt) != GIMPLE_OMP_FOR || gimple_omp_for_kind (ctx->stmt) != GF_OMP_FOR_KIND_FOR) - bad = "#pragma omp for"; + bad = "for"; else if (DECL_FUNCTION_CODE (gimple_call_fndecl (stmt)) == BUILT_IN_GOMP_CANCEL && !integer_zerop (gimple_call_arg (stmt, 1))) @@ -2818,12 +3176,12 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx) if (omp_find_clause (gimple_omp_for_clauses (ctx->stmt), OMP_CLAUSE_NOWAIT)) warning_at (gimple_location (stmt), 0, - "%<#pragma omp cancel for%> inside " + "% inside " "% for construct"); if (omp_find_clause (gimple_omp_for_clauses (ctx->stmt), OMP_CLAUSE_ORDERED)) warning_at (gimple_location (stmt), 0, - "%<#pragma omp cancel for%> inside " + "% inside " "% for construct"); } kind = "for"; @@ -2831,7 +3189,7 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx) case 4: if (gimple_code (ctx->stmt) != GIMPLE_OMP_SECTIONS && gimple_code (ctx->stmt) != GIMPLE_OMP_SECTION) - bad = "#pragma omp sections"; + bad = "sections"; else if (DECL_FUNCTION_CODE (gimple_call_fndecl (stmt)) == BUILT_IN_GOMP_CANCEL && !integer_zerop (gimple_call_arg (stmt, 1))) @@ -2843,7 +3201,7 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx) (ctx->stmt), OMP_CLAUSE_NOWAIT)) warning_at (gimple_location (stmt), 0, - "%<#pragma omp cancel sections%> inside " + "% inside " "% sections construct"); } else @@ -2856,7 +3214,7 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx) (ctx->outer->stmt), OMP_CLAUSE_NOWAIT)) warning_at (gimple_location (stmt), 0, - "%<#pragma omp cancel sections%> inside " + "% inside " "% sections construct"); } } @@ -2867,7 +3225,7 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx) && (!is_taskloop_ctx (ctx) || ctx->outer == NULL || !is_task_ctx (ctx->outer))) - bad = "#pragma omp task"; + bad = "task"; else { for (omp_context *octx = ctx->outer; @@ -2945,14 +3303,14 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx) return true; error_at (gimple_location (stmt), "barrier region may not be closely nested inside " - "of work-sharing, %, %, " - "%, explicit % or % " - "region"); + "of work-sharing, %, %, " + "%, %, explicit % or " + "% region"); return false; } error_at (gimple_location (stmt), "work-sharing region may not be closely nested inside " - "of work-sharing, %, %, " + "of work-sharing, %, %, %, " "%, explicit % or % region"); return false; case GIMPLE_OMP_PARALLEL: @@ -2981,8 +3339,8 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx) case GIMPLE_OMP_TASK: error_at (gimple_location (stmt), "% region may not be closely nested inside " - "of work-sharing, explicit % or % " - "region"); + "of work-sharing, %, explicit % or " + "% region"); return false; case GIMPLE_OMP_PARALLEL: case GIMPLE_OMP_TEAMS: @@ -3200,6 +3558,7 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx) stmt_name = "target exit data"; break; case GF_OMP_TARGET_KIND_OACC_PARALLEL: stmt_name = "parallel"; break; case GF_OMP_TARGET_KIND_OACC_KERNELS: stmt_name = "kernels"; break; + case GF_OMP_TARGET_KIND_OACC_SERIAL: stmt_name = "serial"; break; case GF_OMP_TARGET_KIND_OACC_DATA: stmt_name = "data"; break; case GF_OMP_TARGET_KIND_OACC_UPDATE: stmt_name = "update"; break; case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: @@ -3207,6 +3566,11 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx) case GF_OMP_TARGET_KIND_OACC_DECLARE: stmt_name = "declare"; break; case GF_OMP_TARGET_KIND_OACC_HOST_DATA: stmt_name = "host_data"; break; + case GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_PARALLELIZED: + case GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_GANG_SINGLE: + case GF_OMP_TARGET_KIND_OACC_DATA_KERNELS: + /* OpenACC 'kernels' decomposed parts. */ + stmt_name = "kernels"; break; default: gcc_unreachable (); } switch (gimple_omp_target_kind (ctx->stmt)) @@ -3217,9 +3581,16 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx) ctx_stmt_name = "parallel"; break; case GF_OMP_TARGET_KIND_OACC_KERNELS: ctx_stmt_name = "kernels"; break; + case GF_OMP_TARGET_KIND_OACC_SERIAL: + ctx_stmt_name = "serial"; break; case GF_OMP_TARGET_KIND_OACC_DATA: ctx_stmt_name = "data"; break; case GF_OMP_TARGET_KIND_OACC_HOST_DATA: ctx_stmt_name = "host_data"; break; + case GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_PARALLELIZED: + case GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_GANG_SINGLE: + case GF_OMP_TARGET_KIND_OACC_DATA_KERNELS: + /* OpenACC 'kernels' decomposed parts. */ + ctx_stmt_name = "kernels"; break; default: gcc_unreachable (); } @@ -3321,12 +3692,123 @@ setjmp_or_longjmp_p (const_tree fndecl) return true; tree declname = DECL_NAME (fndecl); - if (!declname) + if (!declname + || (DECL_CONTEXT (fndecl) != NULL_TREE + && TREE_CODE (DECL_CONTEXT (fndecl)) != TRANSLATION_UNIT_DECL) + || !TREE_PUBLIC (fndecl)) return false; + const char *name = IDENTIFIER_POINTER (declname); return !strcmp (name, "setjmp") || !strcmp (name, "longjmp"); } +/* Return true if FNDECL is an omp_* runtime API call. */ + +static bool +omp_runtime_api_call (const_tree fndecl) +{ + tree declname = DECL_NAME (fndecl); + if (!declname + || (DECL_CONTEXT (fndecl) != NULL_TREE + && TREE_CODE (DECL_CONTEXT (fndecl)) != TRANSLATION_UNIT_DECL) + || !TREE_PUBLIC (fndecl)) + return false; + + const char *name = IDENTIFIER_POINTER (declname); + if (strncmp (name, "omp_", 4) != 0) + return false; + + static const char *omp_runtime_apis[] = + { + /* This array has 3 sections. First omp_* calls that don't + have any suffixes. */ + "target_alloc", + "target_associate_ptr", + "target_disassociate_ptr", + "target_free", + "target_is_present", + "target_memcpy", + "target_memcpy_rect", + NULL, + /* Now omp_* calls that are available as omp_* and omp_*_. */ + "capture_affinity", + "destroy_lock", + "destroy_nest_lock", + "display_affinity", + "get_active_level", + "get_affinity_format", + "get_cancellation", + "get_default_device", + "get_dynamic", + "get_initial_device", + "get_level", + "get_max_active_levels", + "get_max_task_priority", + "get_max_threads", + "get_nested", + "get_num_devices", + "get_num_places", + "get_num_procs", + "get_num_teams", + "get_num_threads", + "get_partition_num_places", + "get_place_num", + "get_proc_bind", + "get_team_num", + "get_thread_limit", + "get_thread_num", + "get_wtick", + "get_wtime", + "in_final", + "in_parallel", + "init_lock", + "init_nest_lock", + "is_initial_device", + "pause_resource", + "pause_resource_all", + "set_affinity_format", + "set_lock", + "set_nest_lock", + "test_lock", + "test_nest_lock", + "unset_lock", + "unset_nest_lock", + NULL, + /* And finally calls available as omp_*, omp_*_ and omp_*_8_. */ + "get_ancestor_thread_num", + "get_partition_place_nums", + "get_place_num_procs", + "get_place_proc_ids", + "get_schedule", + "get_team_size", + "set_default_device", + "set_dynamic", + "set_max_active_levels", + "set_nested", + "set_num_threads", + "set_schedule" + }; + + int mode = 0; + for (unsigned i = 0; i < ARRAY_SIZE (omp_runtime_apis); i++) + { + if (omp_runtime_apis[i] == NULL) + { + mode++; + continue; + } + size_t len = strlen (omp_runtime_apis[i]); + if (strncmp (name + 4, omp_runtime_apis[i], len) == 0 + && (name[4 + len] == '\0' + || (mode > 0 + && name[4 + len] == '_' + && (name[4 + len + 1] == '\0' + || (mode > 1 + && strcmp (name + 4 + len + 1, "8_") == 0))))) + return true; + } + return false; +} /* Helper function for scan_omp. @@ -3352,14 +3834,15 @@ scan_omp_1_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p, tree fndecl = gimple_call_fndecl (stmt); if (fndecl) { - if (setjmp_or_longjmp_p (fndecl) - && ctx + if (ctx && gimple_code (ctx->stmt) == GIMPLE_OMP_FOR - && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_SIMD) + && gimple_omp_for_kind (ctx->stmt) == GF_OMP_FOR_KIND_SIMD + && setjmp_or_longjmp_p (fndecl) + && !ctx->loop_p) { remove = true; error_at (gimple_location (stmt), - "setjmp/longjmp inside simd construct"); + "setjmp/longjmp inside % construct"); } else if (DECL_BUILT_IN_CLASS (fndecl) == BUILT_IN_NORMAL) switch (DECL_FUNCTION_CODE (fndecl)) @@ -3376,6 +3859,19 @@ scan_omp_1_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p, default: break; } + else if (ctx) + { + omp_context *octx = ctx; + if (gimple_code (ctx->stmt) == GIMPLE_OMP_SCAN && ctx->outer) + octx = ctx->outer; + if (octx->order_concurrent && omp_runtime_api_call (fndecl)) + { + remove = true; + error_at (gimple_location (stmt), + "OpenMP runtime API call %qD in a region with " + "% clause", fndecl); + } + } } } if (remove) @@ -3417,7 +3913,8 @@ scan_omp_1_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p, if ((gimple_omp_for_kind (as_a (stmt)) == GF_OMP_FOR_KIND_SIMD) && omp_maybe_offloaded_ctx (ctx) - && omp_max_simt_vf ()) + && omp_max_simt_vf () + && gimple_omp_for_collapse (stmt) == 1) scan_omp_simd (gsi, as_a (stmt), ctx); else scan_omp_for (as_a (stmt), ctx); @@ -3444,7 +3941,6 @@ scan_omp_1_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p, case GIMPLE_OMP_MASTER: case GIMPLE_OMP_ORDERED: case GIMPLE_OMP_CRITICAL: - case GIMPLE_OMP_GRID_BODY: ctx = new_omp_context (stmt, ctx); scan_omp (gimple_omp_body_ptr (stmt), ctx); break; @@ -3456,7 +3952,14 @@ scan_omp_1_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p, break; case GIMPLE_OMP_TARGET: - scan_omp_target (as_a (stmt), ctx); + if (is_gimple_omp_offloaded (stmt)) + { + taskreg_nesting_level++; + scan_omp_target (as_a (stmt), ctx); + taskreg_nesting_level--; + } + else + scan_omp_target (as_a (stmt), ctx); break; case GIMPLE_OMP_TEAMS: @@ -3739,11 +4242,8 @@ omp_clause_aligned_alignment (tree clause) /* Otherwise return implementation defined alignment. */ unsigned int al = 1; opt_scalar_mode mode_iter; - auto_vector_sizes sizes; - targetm.vectorize.autovectorize_vector_sizes (&sizes, true); - poly_uint64 vs = 0; - for (unsigned int i = 0; i < sizes.length (); ++i) - vs = ordered_max (vs, sizes[i]); + auto_vector_modes modes; + targetm.vectorize.autovectorize_vector_modes (&modes, true); static enum mode_class classes[] = { MODE_INT, MODE_VECTOR_INT, MODE_FLOAT, MODE_VECTOR_FLOAT }; for (int i = 0; i < 4; i += 2) @@ -3754,17 +4254,16 @@ omp_clause_aligned_alignment (tree clause) machine_mode vmode = targetm.vectorize.preferred_simd_mode (mode); if (GET_MODE_CLASS (vmode) != classes[i + 1]) continue; - while (maybe_ne (vs, 0U) - && known_lt (GET_MODE_SIZE (vmode), vs) - && GET_MODE_2XWIDER_MODE (vmode).exists ()) - vmode = GET_MODE_2XWIDER_MODE (vmode).require (); + machine_mode alt_vmode; + for (unsigned int j = 0; j < modes.length (); ++j) + if (related_vector_mode (modes[j], mode).exists (&alt_vmode) + && known_ge (GET_MODE_SIZE (alt_vmode), GET_MODE_SIZE (vmode))) + vmode = alt_vmode; tree type = lang_hooks.types.type_for_mode (mode, 1); if (type == NULL_TREE || TYPE_MODE (type) != mode) continue; - poly_uint64 nelts = exact_div (GET_MODE_SIZE (vmode), - GET_MODE_SIZE (mode)); - type = build_vector_type (type, nelts); + type = build_vector_type_for_mode (type, vmode); if (TYPE_MODE (type) != vmode) continue; if (TYPE_ALIGN_UNIT (type) > al) @@ -3837,8 +4336,7 @@ lower_rec_simd_input_clauses (tree new_var, omp_context *ctx, DECL_ATTRIBUTES (ivar) = tree_cons (get_identifier ("omp simt private"), NULL, DECL_ATTRIBUTES (ivar)); sctx->simt_eargs.safe_push (build1 (ADDR_EXPR, ptype, ivar)); - tree clobber = build_constructor (type, NULL); - TREE_THIS_VOLATILE (clobber) = 1; + tree clobber = build_clobber (type); gimple *g = gimple_build_assign (ivar, clobber); gimple_seq_add_stmt (&sctx->simt_dlist, g); } @@ -3948,6 +4446,79 @@ task_reduction_read (gimple_seq *ilist, tree tskred_temp, tree type, return v; } +/* Lower early initialization of privatized variable NEW_VAR + if it needs an allocator (has allocate clause). */ + +static bool +lower_private_allocate (tree var, tree new_var, tree &allocator, + tree &allocate_ptr, gimple_seq *ilist, + omp_context *ctx, bool is_ref, tree size) +{ + if (allocator) + return false; + gcc_assert (allocate_ptr == NULL_TREE); + if (ctx->allocate_map + && (DECL_P (new_var) || (TYPE_P (new_var) && size))) + if (tree *allocatorp = ctx->allocate_map->get (var)) + allocator = *allocatorp; + if (allocator == NULL_TREE) + return false; + if (!is_ref && omp_is_reference (var)) + { + allocator = NULL_TREE; + return false; + } + + if (TREE_CODE (allocator) != INTEGER_CST) + allocator = build_outer_var_ref (allocator, ctx); + allocator = fold_convert (pointer_sized_int_node, allocator); + if (TREE_CODE (allocator) != INTEGER_CST) + { + tree var = create_tmp_var (TREE_TYPE (allocator)); + gimplify_assign (var, allocator, ilist); + allocator = var; + } + + tree ptr_type, align, sz = size; + if (TYPE_P (new_var)) + { + ptr_type = build_pointer_type (new_var); + align = build_int_cst (size_type_node, TYPE_ALIGN_UNIT (new_var)); + } + else if (is_ref) + { + ptr_type = build_pointer_type (TREE_TYPE (TREE_TYPE (new_var))); + align = build_int_cst (size_type_node, + TYPE_ALIGN_UNIT (TREE_TYPE (ptr_type))); + } + else + { + ptr_type = build_pointer_type (TREE_TYPE (new_var)); + align = build_int_cst (size_type_node, DECL_ALIGN_UNIT (new_var)); + if (sz == NULL_TREE) + sz = fold_convert (size_type_node, DECL_SIZE_UNIT (new_var)); + } + if (TREE_CODE (sz) != INTEGER_CST) + { + tree szvar = create_tmp_var (size_type_node); + gimplify_assign (szvar, sz, ilist); + sz = szvar; + } + allocate_ptr = create_tmp_var (ptr_type); + tree a = builtin_decl_explicit (BUILT_IN_GOMP_ALLOC); + gimple *g = gimple_build_call (a, 3, align, sz, allocator); + gimple_call_set_lhs (g, allocate_ptr); + gimple_seq_add_stmt (ilist, g); + if (!is_ref) + { + tree x = build_simple_mem_ref (allocate_ptr); + TREE_THIS_NOTRAP (x) = 1; + SET_DECL_VALUE_EXPR (new_var, x); + DECL_HAS_VALUE_EXPR_P (new_var) = 1; + } + return true; +} + /* Generate code to implement the input clauses, FIRSTPRIVATE and COPYIN, from the receiver (aka child) side and initializers for REFERENCE_TYPE private variables. Initialization statements go in ILIST, while calls @@ -3963,7 +4534,7 @@ lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist, bool reduction_omp_orig_ref = false; int pass; bool is_simd = (gimple_code (ctx->stmt) == GIMPLE_OMP_FOR - && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_SIMD); + && gimple_omp_for_kind (ctx->stmt) == GF_OMP_FOR_KIND_SIMD); omplow_simd_context sctx = omplow_simd_context (); tree simt_lane = NULL_TREE, simtrec = NULL_TREE; tree ivar = NULL_TREE, lvar = NULL_TREE, uid = NULL_TREE; @@ -4113,6 +4684,7 @@ lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist, bool task_reduction_p = false; bool task_reduction_needs_orig_p = false; tree cond = NULL_TREE; + tree allocator, allocate_ptr; switch (c_kind) { @@ -4249,6 +4821,8 @@ lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist, if (task_reduction_p != (pass >= 2)) continue; + allocator = NULL_TREE; + allocate_ptr = NULL_TREE; new_var = var = OMP_CLAUSE_DECL (c); if ((c_kind == OMP_CLAUSE_REDUCTION || c_kind == OMP_CLAUSE_IN_REDUCTION) @@ -4357,7 +4931,23 @@ lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist, tree type = TREE_TYPE (d); gcc_assert (TREE_CODE (type) == ARRAY_TYPE); tree v = TYPE_MAX_VALUE (TYPE_DOMAIN (type)); + tree sz = v; const char *name = get_name (orig_var); + if (pass != 3 && !TREE_CONSTANT (v)) + { + tree t = maybe_lookup_decl (v, ctx); + if (t) + v = t; + else + v = maybe_lookup_decl_in_outer_ctx (v, ctx); + gimplify_expr (&v, ilist, NULL, is_gimple_val, fb_rvalue); + t = fold_build2_loc (clause_loc, PLUS_EXPR, + TREE_TYPE (v), v, + build_int_cst (TREE_TYPE (v), 1)); + sz = fold_build2_loc (clause_loc, MULT_EXPR, + TREE_TYPE (v), t, + TYPE_SIZE_UNIT (TREE_TYPE (type))); + } if (pass == 3) { tree xv = create_tmp_var (ptr_type_node); @@ -4415,6 +5005,13 @@ lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist, gimplify_assign (cond, x, ilist); x = xv; } + else if (lower_private_allocate (var, type, allocator, + allocate_ptr, ilist, ctx, + true, + TREE_CONSTANT (v) + ? TYPE_SIZE_UNIT (type) + : sz)) + x = allocate_ptr; else if (TREE_CONSTANT (v)) { x = create_tmp_var_raw (type, name); @@ -4426,20 +5023,8 @@ lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist, { tree atmp = builtin_decl_explicit (BUILT_IN_ALLOCA_WITH_ALIGN); - tree t = maybe_lookup_decl (v, ctx); - if (t) - v = t; - else - v = maybe_lookup_decl_in_outer_ctx (v, ctx); - gimplify_expr (&v, ilist, NULL, is_gimple_val, fb_rvalue); - t = fold_build2_loc (clause_loc, PLUS_EXPR, - TREE_TYPE (v), v, - build_int_cst (TREE_TYPE (v), 1)); - t = fold_build2_loc (clause_loc, MULT_EXPR, - TREE_TYPE (v), t, - TYPE_SIZE_UNIT (TREE_TYPE (type))); tree al = size_int (TYPE_ALIGN (TREE_TYPE (type))); - x = build_call_expr_loc (clause_loc, atmp, 2, t, al); + x = build_call_expr_loc (clause_loc, atmp, 2, sz, al); } tree ptype = build_pointer_type (TREE_TYPE (type)); @@ -4701,6 +5286,12 @@ lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist, gimple_seq_add_stmt (dlist, g); gimple_seq_add_stmt (dlist, gimple_build_label (end2)); } + if (allocator) + { + tree f = builtin_decl_explicit (BUILT_IN_GOMP_FREE); + g = gimple_build_call (f, 2, allocate_ptr, allocator); + gimple_seq_add_stmt (dlist, g); + } continue; } else if (pass == 2) @@ -4785,8 +5376,7 @@ lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist, if (c_kind != OMP_CLAUSE_FIRSTPRIVATE || !is_task_ctx (ctx)) { - gcall *stmt; - tree tmp, atmp; + tree tmp; ptr = DECL_VALUE_EXPR (new_var); gcc_assert (TREE_CODE (ptr) == INDIRECT_REF); @@ -4794,15 +5384,25 @@ lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist, gcc_assert (DECL_P (ptr)); x = TYPE_SIZE_UNIT (TREE_TYPE (new_var)); - /* void *tmp = __builtin_alloca */ - atmp = builtin_decl_explicit (BUILT_IN_ALLOCA_WITH_ALIGN); - stmt = gimple_build_call (atmp, 2, x, - size_int (DECL_ALIGN (var))); - tmp = create_tmp_var_raw (ptr_type_node); - gimple_add_tmp_var (tmp); - gimple_call_set_lhs (stmt, tmp); - - gimple_seq_add_stmt (ilist, stmt); + if (lower_private_allocate (var, new_var, allocator, + allocate_ptr, ilist, ctx, + false, x)) + tmp = allocate_ptr; + else + { + /* void *tmp = __builtin_alloca */ + tree atmp + = builtin_decl_explicit (BUILT_IN_ALLOCA_WITH_ALIGN); + gcall *stmt + = gimple_build_call (atmp, 2, x, + size_int (DECL_ALIGN (var))); + cfun->calls_alloca = 1; + tmp = create_tmp_var_raw (ptr_type_node); + gimple_add_tmp_var (tmp); + gimple_call_set_lhs (stmt, tmp); + + gimple_seq_add_stmt (ilist, stmt); + } x = fold_convert_loc (clause_loc, TREE_TYPE (ptr), tmp); gimplify_assign (ptr, x, ilist); @@ -4824,8 +5424,23 @@ lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist, if (c_kind == OMP_CLAUSE_FIRSTPRIVATE && is_task_ctx (ctx)) { x = build_receiver_ref (var, false, ctx); - x = build_fold_addr_expr_loc (clause_loc, x); + if (ctx->allocate_map) + if (tree *allocatep = ctx->allocate_map->get (var)) + { + allocator = *allocatep; + if (TREE_CODE (allocator) != INTEGER_CST) + allocator = build_outer_var_ref (allocator, ctx); + allocator = fold_convert (pointer_sized_int_node, + allocator); + allocate_ptr = unshare_expr (x); + } + if (allocator == NULL_TREE) + x = build_fold_addr_expr_loc (clause_loc, x); } + else if (lower_private_allocate (var, new_var, allocator, + allocate_ptr, + ilist, ctx, true, x)) + x = allocate_ptr; else if (TREE_CONSTANT (x)) { /* For reduction in SIMD loop, defer adding the @@ -4936,13 +5551,27 @@ lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist, x = NULL; do_private: tree nx; - nx = lang_hooks.decls.omp_clause_default_ctor - (c, unshare_expr (new_var), x); + bool copy_ctor; + copy_ctor = false; + lower_private_allocate (var, new_var, allocator, allocate_ptr, + ilist, ctx, false, NULL_TREE); + nx = unshare_expr (new_var); + if (is_simd + && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LASTPRIVATE + && OMP_CLAUSE_LASTPRIVATE_LOOP_IV (c)) + copy_ctor = true; + if (copy_ctor) + nx = lang_hooks.decls.omp_clause_copy_ctor (c, nx, x); + else + nx = lang_hooks.decls.omp_clause_default_ctor (c, nx, x); if (is_simd) { tree y = lang_hooks.decls.omp_clause_dtor (c, new_var); if ((TREE_ADDRESSABLE (new_var) || nx || y - || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LASTPRIVATE + || (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LASTPRIVATE + && (gimple_omp_for_collapse (ctx->stmt) != 1 + || (gimple_omp_for_index (ctx->stmt, 0) + != new_var))) || OMP_CLAUSE_CODE (c) == OMP_CLAUSE__CONDTEMP_ || omp_is_reference (var)) && lower_rec_simd_input_clauses (new_var, ctx, &sctx, @@ -4959,8 +5588,16 @@ lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist, } if (nx) - x = lang_hooks.decls.omp_clause_default_ctor - (c, unshare_expr (ivar), x); + { + tree iv = unshare_expr (ivar); + if (copy_ctor) + x = lang_hooks.decls.omp_clause_copy_ctor (c, iv, + x); + else + x = lang_hooks.decls.omp_clause_default_ctor (c, + iv, + x); + } else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE__CONDTEMP_) { x = build2 (MODIFY_EXPR, TREE_TYPE (ivar), @@ -5063,6 +5700,25 @@ lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist, x = lang_hooks.decls.omp_clause_dtor (c, new_var); if (x) gimplify_and_add (x, dlist); + if (allocator) + { + if (!is_gimple_val (allocator)) + { + tree avar = create_tmp_var (TREE_TYPE (allocator)); + gimplify_assign (avar, allocator, dlist); + allocator = avar; + } + if (!is_gimple_val (allocate_ptr)) + { + tree apvar = create_tmp_var (TREE_TYPE (allocate_ptr)); + gimplify_assign (apvar, allocate_ptr, dlist); + allocate_ptr = apvar; + } + tree f = builtin_decl_explicit (BUILT_IN_GOMP_FREE); + gimple *g + = gimple_build_call (f, 2, allocate_ptr, allocator); + gimple_seq_add_stmt (dlist, g); + } break; case OMP_CLAUSE_LINEAR: @@ -5086,6 +5742,18 @@ lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist, || use_pointer_for_field (var, NULL)) { x = build_receiver_ref (var, false, ctx); + if (ctx->allocate_map) + if (tree *allocatep = ctx->allocate_map->get (var)) + { + allocator = *allocatep; + if (TREE_CODE (allocator) != INTEGER_CST) + allocator = build_outer_var_ref (allocator, ctx); + allocator = fold_convert (pointer_sized_int_node, + allocator); + allocate_ptr = unshare_expr (x); + x = build_simple_mem_ref (x); + TREE_THIS_NOTRAP (x) = 1; + } SET_DECL_VALUE_EXPR (new_var, x); DECL_HAS_VALUE_EXPR_P (new_var) = 1; goto do_dtor; @@ -5104,6 +5772,8 @@ lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist, goto do_dtor; } do_firstprivate: + lower_private_allocate (var, new_var, allocator, allocate_ptr, + ilist, ctx, false, NULL_TREE); x = build_outer_var_ref (var, ctx); if (is_simd) { @@ -5291,6 +5961,9 @@ lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist, } else { + lower_private_allocate (var, new_var, allocator, + allocate_ptr, ilist, ctx, false, + NULL_TREE); x = build_outer_var_ref (var, ctx); if (omp_is_reference (var) @@ -5687,6 +6360,9 @@ lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist, } else { + lower_private_allocate (var, new_var, allocator, + allocate_ptr, ilist, ctx, + false, NULL_TREE); if (omp_is_reference (var) && is_simd) handle_simd_reference (clause_loc, new_vard, ilist); if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION @@ -5701,6 +6377,8 @@ lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist, ref = build_outer_var_ref (var, ctx); gimplify_assign (ref, x, dlist); } + if (allocator) + goto do_dtor; } } break; @@ -5712,8 +6390,7 @@ lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist, } if (tskred_avar) { - tree clobber = build_constructor (TREE_TYPE (tskred_avar), NULL); - TREE_THIS_VOLATILE (clobber) = 1; + tree clobber = build_clobber (TREE_TYPE (tskred_avar)); gimple_seq_add_stmt (ilist, gimple_build_assign (tskred_avar, clobber)); } @@ -5938,7 +6615,7 @@ lower_lastprivate_conditional_clauses (tree *clauses, omp_context *ctx) tree cond_ptr = NULL_TREE; tree iter_var = NULL_TREE; bool is_simd = (gimple_code (ctx->stmt) == GIMPLE_OMP_FOR - && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_SIMD); + && gimple_omp_for_kind (ctx->stmt) == GF_OMP_FOR_KIND_SIMD); tree next = *clauses; for (tree c = *clauses; c; c = OMP_CLAUSE_CHAIN (c)) if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LASTPRIVATE @@ -6070,7 +6747,7 @@ lower_lastprivate_clauses (tree clauses, tree predicate, gimple_seq *body_p, bool maybe_simt = false; if (gimple_code (ctx->stmt) == GIMPLE_OMP_FOR - && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_SIMD) + && gimple_omp_for_kind (ctx->stmt) == GF_OMP_FOR_KIND_SIMD) { maybe_simt = omp_find_clause (orig_clauses, OMP_CLAUSE__SIMT_); simduid = omp_find_clause (orig_clauses, OMP_CLAUSE__SIMDUID_); @@ -6263,9 +6940,9 @@ lower_lastprivate_clauses (tree clauses, tree predicate, gimple_seq *body_p, x = NULL_TREE; if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LASTPRIVATE - && OMP_CLAUSE_LASTPRIVATE_TASKLOOP_IV (c)) + && OMP_CLAUSE_LASTPRIVATE_LOOP_IV (c) + && is_taskloop_ctx (ctx)) { - gcc_checking_assert (is_taskloop_ctx (ctx)); tree ovar = maybe_lookup_decl_in_outer_ctx (var, ctx->outer->outer); if (is_global_var (ovar)) @@ -6338,6 +7015,11 @@ lower_oacc_reductions (location_t loc, tree clauses, tree level, bool inner, for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION) { + /* No 'reduction' clauses on OpenACC 'kernels'. */ + gcc_checking_assert (!is_oacc_kernels (ctx)); + /* Likewise, on OpenACC 'kernels' decomposed parts. */ + gcc_checking_assert (!is_oacc_kernels_decomposed_part (ctx)); + tree orig = OMP_CLAUSE_DECL (c); tree var = maybe_lookup_decl (orig, ctx); tree ref_to_res = NULL_TREE; @@ -6375,8 +7057,11 @@ lower_oacc_reductions (location_t loc, tree clauses, tree level, bool inner, break; case GIMPLE_OMP_TARGET: - if (gimple_omp_target_kind (probe->stmt) - != GF_OMP_TARGET_KIND_OACC_PARALLEL) + /* No 'reduction' clauses inside OpenACC 'kernels' + regions. */ + gcc_checking_assert (!is_oacc_kernels (probe)); + + if (!is_gimple_omp_offloaded (probe->stmt)) goto do_lookup; cls = gimple_omp_target_clauses (probe->stmt); @@ -6552,7 +7237,7 @@ lower_reduction_clauses (tree clauses, gimple_seq *stmt_seqp, /* SIMD reductions are handled in lower_rec_input_clauses. */ if (gimple_code (ctx->stmt) == GIMPLE_OMP_FOR - && gimple_omp_for_kind (ctx->stmt) & GF_OMP_FOR_SIMD) + && gimple_omp_for_kind (ctx->stmt) == GF_OMP_FOR_KIND_SIMD) return; /* inscan reductions are handled elsewhere. */ @@ -7056,7 +7741,10 @@ lower_send_shared_vars (gimple_seq *ilist, gimple_seq *olist, omp_context *ctx) continue; nvar = maybe_lookup_decl (ovar, ctx); - if (!nvar || !DECL_HAS_VALUE_EXPR_P (nvar)) + if (!nvar + || !DECL_HAS_VALUE_EXPR_P (nvar) + || (ctx->allocate_map + && ctx->allocate_map->get (ovar))) continue; /* If CTX is a nested parallel directive. Find the immediately @@ -7180,11 +7868,29 @@ lower_oacc_head_mark (location_t loc, tree ddvar, tree clauses, tag |= OLF_GANG_STATIC; } - /* In a parallel region, loops are implicitly INDEPENDENT. */ omp_context *tgt = enclosing_target_ctx (ctx); - if (!tgt || is_oacc_parallel (tgt)) + if (!tgt || is_oacc_parallel_or_serial (tgt)) + ; + else if (is_oacc_kernels (tgt)) + /* Not using this loops handling inside OpenACC 'kernels' regions. */ + gcc_unreachable (); + else if (is_oacc_kernels_decomposed_part (tgt)) + ; + else + gcc_unreachable (); + + /* In a parallel region, loops are implicitly INDEPENDENT. */ + if (!tgt || is_oacc_parallel_or_serial (tgt)) tag |= OLF_INDEPENDENT; + /* Loops inside OpenACC 'kernels' decomposed parts' regions are expected to + have an explicit 'seq' or 'independent' clause, and no 'auto' clause. */ + if (tgt && is_oacc_kernels_decomposed_part (tgt)) + { + gcc_assert (tag & (OLF_SEQ | OLF_INDEPENDENT)); + gcc_assert (!(tag & OLF_AUTO)); + } + if (tag & OLF_TILE) /* Tiling could use all 3 levels. */ levels = 3; @@ -7669,8 +8375,7 @@ lower_omp_single (gimple_stmt_iterator *gsi_p, omp_context *ctx) if (ctx->record_type) { gimple_stmt_iterator gsi = gsi_start (bind_body_tail); - tree clobber = build_constructor (ctx->record_type, NULL); - TREE_THIS_VOLATILE (clobber) = 1; + tree clobber = build_clobber (ctx->record_type); gsi_insert_after (&gsi, gimple_build_assign (ctx->sender_decl, clobber), GSI_SAME_STMT); } @@ -8768,7 +9473,7 @@ lower_omp_scan (gimple_stmt_iterator *gsi_p, omp_context *ctx) bool input_phase = has_clauses ^ octx->scan_inclusive; bool is_simd = (gimple_code (octx->stmt) == GIMPLE_OMP_FOR - && (gimple_omp_for_kind (octx->stmt) & GF_OMP_FOR_SIMD)); + && gimple_omp_for_kind (octx->stmt) == GF_OMP_FOR_KIND_SIMD); bool is_for = (gimple_code (octx->stmt) == GIMPLE_OMP_FOR && gimple_omp_for_kind (octx->stmt) == GF_OMP_FOR_KIND_FOR && !gimple_omp_for_combined_p (octx->stmt)); @@ -9157,65 +9862,65 @@ lower_omp_for_lastprivate (struct omp_for_data *fd, gimple_seq *body_p, cond_code = EQ_EXPR; } - if (gimple_omp_for_kind (fd->for_stmt) == GF_OMP_FOR_KIND_GRID_LOOP - || gimple_omp_for_grid_phony (fd->for_stmt)) - cond = omp_grid_lastprivate_predicate (fd); - else + tree n2 = fd->loop.n2; + if (fd->collapse > 1 + && TREE_CODE (n2) != INTEGER_CST + && gimple_omp_for_combined_into_p (fd->for_stmt)) { - tree n2 = fd->loop.n2; - if (fd->collapse > 1 - && TREE_CODE (n2) != INTEGER_CST - && gimple_omp_for_combined_into_p (fd->for_stmt)) + struct omp_context *taskreg_ctx = NULL; + if (gimple_code (ctx->outer->stmt) == GIMPLE_OMP_FOR) { - struct omp_context *taskreg_ctx = NULL; - if (gimple_code (ctx->outer->stmt) == GIMPLE_OMP_FOR) + gomp_for *gfor = as_a (ctx->outer->stmt); + if (gimple_omp_for_kind (gfor) == GF_OMP_FOR_KIND_FOR + || gimple_omp_for_kind (gfor) == GF_OMP_FOR_KIND_DISTRIBUTE) { - gomp_for *gfor = as_a (ctx->outer->stmt); - if (gimple_omp_for_kind (gfor) == GF_OMP_FOR_KIND_FOR - || gimple_omp_for_kind (gfor) == GF_OMP_FOR_KIND_DISTRIBUTE) + if (gimple_omp_for_combined_into_p (gfor)) { - 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; - omp_extract_for_data (gfor, &outer_fd, NULL); - n2 = fold_convert (TREE_TYPE (n2), outer_fd.loop.n2); - } + gcc_assert (ctx->outer->outer + && is_parallel_ctx (ctx->outer->outer)); + taskreg_ctx = ctx->outer->outer; } - else if (gimple_omp_for_kind (gfor) == GF_OMP_FOR_KIND_TASKLOOP) - taskreg_ctx = ctx->outer->outer; - } - else if (is_taskreg_ctx (ctx->outer)) - taskreg_ctx = ctx->outer; - if (taskreg_ctx) - { - int i; - tree taskreg_clauses - = gimple_omp_taskreg_clauses (taskreg_ctx->stmt); - tree innerc = omp_find_clause (taskreg_clauses, - OMP_CLAUSE__LOOPTEMP_); - gcc_assert (innerc); - for (i = 0; i < fd->collapse; i++) + else { - innerc = omp_find_clause (OMP_CLAUSE_CHAIN (innerc), - OMP_CLAUSE__LOOPTEMP_); - gcc_assert (innerc); + struct omp_for_data outer_fd; + omp_extract_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) + taskreg_ctx = ctx->outer->outer; + } + else if (is_taskreg_ctx (ctx->outer)) + taskreg_ctx = ctx->outer; + if (taskreg_ctx) + { + int i; + tree taskreg_clauses + = gimple_omp_taskreg_clauses (taskreg_ctx->stmt); + tree innerc = omp_find_clause (taskreg_clauses, + OMP_CLAUSE__LOOPTEMP_); + gcc_assert (innerc); + int count = fd->collapse; + if (fd->non_rect + && fd->last_nonrect == fd->first_nonrect + 1) + if (tree v = gimple_omp_for_index (fd->for_stmt, fd->last_nonrect)) + if (!TYPE_UNSIGNED (TREE_TYPE (v))) + count += 4; + for (i = 0; i < count; i++) + { innerc = omp_find_clause (OMP_CLAUSE_CHAIN (innerc), OMP_CLAUSE__LOOPTEMP_); - if (innerc) - n2 = fold_convert (TREE_TYPE (n2), - lookup_decl (OMP_CLAUSE_DECL (innerc), - taskreg_ctx)); + gcc_assert (innerc); } + innerc = omp_find_clause (OMP_CLAUSE_CHAIN (innerc), + OMP_CLAUSE__LOOPTEMP_); + if (innerc) + n2 = fold_convert (TREE_TYPE (n2), + lookup_decl (OMP_CLAUSE_DECL (innerc), + taskreg_ctx)); } - cond = build2 (cond_code, boolean_type_node, fd->loop.v, n2); } + cond = build2 (cond_code, boolean_type_node, fd->loop.v, n2); clauses = gimple_omp_for_clauses (fd->for_stmt); stmts = NULL; @@ -9254,7 +9959,7 @@ omp_find_scan (gimple_stmt_iterator *gsi_p, bool *handled_ops_p, WALK_SUBSTMTS; case GIMPLE_OMP_FOR: - if ((gimple_omp_for_kind (stmt) & GF_OMP_FOR_SIMD) + if (gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_SIMD && gimple_omp_for_combined_into_p (stmt)) *handled_ops_p = false; break; @@ -9961,7 +10666,6 @@ lower_omp_for_scan (gimple_seq *body_p, gimple_seq *dlist, gomp_for *stmt, gimple_seq_add_stmt (body_p, g); tree cplx = create_tmp_var (build_complex_type (unsigned_type_node, false)); - DECL_GIMPLE_REG_P (cplx) = 1; g = gimple_build_call_internal (IFN_MUL_OVERFLOW, 2, thread_nump1, twok); gimple_call_set_lhs (g, cplx); gimple_seq_add_stmt (body_p, g); @@ -10125,12 +10829,26 @@ lower_omp_for (gimple_stmt_iterator *gsi_p, omp_context *ctx) if (fd.collapse > 1 && TREE_CODE (fd.loop.n2) != INTEGER_CST) count += fd.collapse - 1; + size_t count2 = 0; + tree type2 = NULL_TREE; bool taskreg_for = (gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_FOR || gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_TASKLOOP); tree outerc = NULL, *pc = gimple_omp_for_clauses_ptr (stmt); tree simtc = NULL; tree clauses = *pc; + if (fd.collapse > 1 + && fd.non_rect + && fd.last_nonrect == fd.first_nonrect + 1 + && TREE_CODE (fd.loop.n2) != INTEGER_CST) + if (tree v = gimple_omp_for_index (stmt, fd.last_nonrect)) + if (!TYPE_UNSIGNED (TREE_TYPE (v))) + { + v = gimple_omp_for_index (stmt, fd.first_nonrect); + type2 = TREE_TYPE (v); + count++; + count2 = 3; + } if (taskreg_for) outerc = omp_find_clause (gimple_omp_taskreg_clauses (ctx->outer->stmt), @@ -10138,7 +10856,7 @@ lower_omp_for (gimple_stmt_iterator *gsi_p, omp_context *ctx) if (ctx->simt_stmt) simtc = omp_find_clause (gimple_omp_for_clauses (ctx->simt_stmt), OMP_CLAUSE__LOOPTEMP_); - for (i = 0; i < count; i++) + for (i = 0; i < count + count2; i++) { tree temp; if (taskreg_for) @@ -10157,7 +10875,7 @@ lower_omp_for (gimple_stmt_iterator *gsi_p, omp_context *ctx) if (ctx->simt_stmt) temp = OMP_CLAUSE_DECL (simtc); else - temp = create_tmp_var (type); + temp = create_tmp_var (i >= count ? type2 : type); insert_decl_map (&ctx->outer->cb, temp, temp); } *pc = build_omp_clause (UNKNOWN_LOCATION, OMP_CLAUSE__LOOPTEMP_); @@ -10213,13 +10931,31 @@ lower_omp_for (gimple_stmt_iterator *gsi_p, omp_context *ctx) for (i = 0; i < gimple_omp_for_collapse (stmt); i++) { rhs_p = gimple_omp_for_initial_ptr (stmt, i); - if (!is_gimple_min_invariant (*rhs_p)) + if (TREE_CODE (*rhs_p) == TREE_VEC) + { + if (!is_gimple_min_invariant (TREE_VEC_ELT (*rhs_p, 1))) + TREE_VEC_ELT (*rhs_p, 1) + = get_formal_tmp_var (TREE_VEC_ELT (*rhs_p, 1), &cnt_list); + if (!is_gimple_min_invariant (TREE_VEC_ELT (*rhs_p, 2))) + TREE_VEC_ELT (*rhs_p, 2) + = get_formal_tmp_var (TREE_VEC_ELT (*rhs_p, 2), &cnt_list); + } + else if (!is_gimple_min_invariant (*rhs_p)) *rhs_p = get_formal_tmp_var (*rhs_p, &cnt_list); else if (TREE_CODE (*rhs_p) == ADDR_EXPR) recompute_tree_invariant_for_addr_expr (*rhs_p); rhs_p = gimple_omp_for_final_ptr (stmt, i); - if (!is_gimple_min_invariant (*rhs_p)) + if (TREE_CODE (*rhs_p) == TREE_VEC) + { + if (!is_gimple_min_invariant (TREE_VEC_ELT (*rhs_p, 1))) + TREE_VEC_ELT (*rhs_p, 1) + = get_formal_tmp_var (TREE_VEC_ELT (*rhs_p, 1), &cnt_list); + if (!is_gimple_min_invariant (TREE_VEC_ELT (*rhs_p, 2))) + TREE_VEC_ELT (*rhs_p, 2) + = get_formal_tmp_var (TREE_VEC_ELT (*rhs_p, 2), &cnt_list); + } + else if (!is_gimple_min_invariant (*rhs_p)) *rhs_p = get_formal_tmp_var (*rhs_p, &cnt_list); else if (TREE_CODE (*rhs_p) == ADDR_EXPR) recompute_tree_invariant_for_addr_expr (*rhs_p); @@ -10260,24 +10996,17 @@ lower_omp_for (gimple_stmt_iterator *gsi_p, omp_context *ctx) ctx); } - bool phony_loop = (gimple_omp_for_kind (stmt) != GF_OMP_FOR_KIND_GRID_LOOP - && gimple_omp_for_grid_phony (stmt)); if ((ctx->scan_inclusive || ctx->scan_exclusive) && gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_FOR) - { - gcc_assert (!phony_loop); - lower_omp_for_scan (&body, &dlist, stmt, &fd, ctx); - } + lower_omp_for_scan (&body, &dlist, stmt, &fd, ctx); else { - if (!phony_loop) - gimple_seq_add_stmt (&body, stmt); + gimple_seq_add_stmt (&body, stmt); gimple_seq_add_seq (&body, gimple_omp_body (stmt)); } - if (!phony_loop) - gimple_seq_add_stmt (&body, gimple_build_omp_continue (fd.loop.v, - fd.loop.v)); + gimple_seq_add_stmt (&body, gimple_build_omp_continue (fd.loop.v, + fd.loop.v)); /* After the loop, add exit clauses. */ lower_reduction_clauses (gimple_omp_for_clauses (stmt), &body, &clist, ctx); @@ -10306,19 +11035,16 @@ lower_omp_for (gimple_stmt_iterator *gsi_p, omp_context *ctx) body = maybe_catch_exception (body); - if (!phony_loop) - { - /* Region exit marker goes at the end of the loop body. */ - gimple *g = gimple_build_omp_return (fd.have_nowait); - gimple_seq_add_stmt (&body, g); + /* Region exit marker goes at the end of the loop body. */ + gimple *g = gimple_build_omp_return (fd.have_nowait); + gimple_seq_add_stmt (&body, g); - gimple_seq_add_seq (&body, tred_dlist); + gimple_seq_add_seq (&body, tred_dlist); - maybe_add_implicit_barrier_cancel (ctx, g, &body); + maybe_add_implicit_barrier_cancel (ctx, g, &body); - if (rclauses) - OMP_CLAUSE_DECL (rclauses) = rtmp; - } + if (rclauses) + OMP_CLAUSE_DECL (rclauses) = rtmp; /* Add OpenACC joining and reduction markers just after the loop. */ if (oacc_tail) @@ -10614,7 +11340,35 @@ create_task_copyfn (gomp_task *task_stmt, omp_context *ctx) if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_FIRSTPRIVATE) t = build2 (MODIFY_EXPR, TREE_TYPE (dst), dst, src); else - t = lang_hooks.decls.omp_clause_copy_ctor (c, dst, src); + { + if (ctx->allocate_map) + if (tree *allocatorp = ctx->allocate_map->get (decl)) + { + tree allocator = *allocatorp; + if (TREE_CODE (allocator) != INTEGER_CST) + { + n = splay_tree_lookup (ctx->sfield_map, + (splay_tree_key) allocator); + allocator = (tree) n->value; + if (tcctx.cb.decl_map) + allocator = *tcctx.cb.decl_map->get (allocator); + tree a = build_simple_mem_ref_loc (loc, sarg); + allocator = omp_build_component_ref (a, allocator); + } + allocator = fold_convert (pointer_sized_int_node, allocator); + tree a = builtin_decl_explicit (BUILT_IN_GOMP_ALLOC); + tree align = build_int_cst (size_type_node, + DECL_ALIGN_UNIT (decl)); + tree sz = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (dst))); + tree ptr = build_call_expr_loc (loc, a, 3, align, sz, + allocator); + ptr = fold_convert (TREE_TYPE (dst), ptr); + t = build2 (MODIFY_EXPR, TREE_TYPE (dst), dst, ptr); + append_to_statement_list (t, &list); + dst = build_simple_mem_ref_loc (loc, dst); + } + t = lang_hooks.decls.omp_clause_copy_ctor (c, dst, src); + } append_to_statement_list (t, &list); break; case OMP_CLAUSE_PRIVATE: @@ -10802,8 +11556,7 @@ lower_depend_clauses (tree *pclauses, gimple_seq *iseq, gimple_seq *oseq) OMP_CLAUSE_DECL (c) = build_fold_addr_expr (array); OMP_CLAUSE_CHAIN (c) = *pclauses; *pclauses = c; - tree clobber = build_constructor (type, NULL); - TREE_THIS_VOLATILE (clobber) = 1; + tree clobber = build_clobber (type); g = gimple_build_assign (array, clobber); gimple_seq_add_stmt (oseq, g); } @@ -10902,17 +11655,9 @@ lower_omp_taskreg (gimple_stmt_iterator *gsi_p, omp_context *ctx) gimple_seq par_olist = NULL; gimple_seq par_ilist = NULL; gimple_seq par_rlist = NULL; - bool phony_construct = gimple_code (stmt) == GIMPLE_OMP_PARALLEL - && gimple_omp_parallel_grid_phony (as_a (stmt)); - if (phony_construct && ctx->record_type) - { - gcc_checking_assert (!ctx->receiver_decl); - ctx->receiver_decl = create_tmp_var - (build_reference_type (ctx->record_type), ".omp_rec"); - } lower_rec_input_clauses (clauses, &par_ilist, &par_olist, ctx, NULL); lower_omp (&par_body, ctx); - if (gimple_code (stmt) == GIMPLE_OMP_PARALLEL) + if (gimple_code (stmt) != GIMPLE_OMP_TASK) lower_reduction_clauses (clauses, &par_rlist, NULL, ctx); /* Declare all the variables created by mapping and the variables @@ -10938,8 +11683,7 @@ lower_omp_taskreg (gimple_stmt_iterator *gsi_p, omp_context *ctx) if (ctx->record_type) { - tree clobber = build_constructor (TREE_TYPE (ctx->sender_decl), NULL); - TREE_THIS_VOLATILE (clobber) = 1; + tree clobber = build_clobber (TREE_TYPE (ctx->sender_decl)); gimple_seq_add_stmt (&olist, gimple_build_assign (ctx->sender_decl, clobber)); } @@ -10969,11 +11713,8 @@ lower_omp_taskreg (gimple_stmt_iterator *gsi_p, omp_context *ctx) gimple_seq_add_stmt (&new_body, gimple_build_omp_continue (integer_zero_node, integer_zero_node)); - if (!phony_construct) - { - gimple_seq_add_stmt (&new_body, gimple_build_omp_return (false)); - gimple_omp_set_body (stmt, new_body); - } + gimple_seq_add_stmt (&new_body, gimple_build_omp_return (false)); + gimple_omp_set_body (stmt, new_body); if (dep_bind && gimple_bind_block (par_bind) == NULL_TREE) bind = gimple_build_bind (NULL, NULL, make_node (BLOCK)); @@ -10981,10 +11722,7 @@ lower_omp_taskreg (gimple_stmt_iterator *gsi_p, omp_context *ctx) bind = gimple_build_bind (NULL, NULL, gimple_bind_block (par_bind)); gsi_replace (gsi_p, dep_bind ? dep_bind : bind, true); gimple_bind_add_seq (bind, ilist); - if (!phony_construct) - gimple_bind_add_stmt (bind, stmt); - else - gimple_bind_add_seq (bind, new_body); + gimple_bind_add_stmt (bind, stmt); gimple_bind_add_seq (bind, olist); pop_gimplify_context (NULL); @@ -11024,14 +11762,18 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) case GF_OMP_TARGET_KIND_EXIT_DATA: case GF_OMP_TARGET_KIND_OACC_PARALLEL: case GF_OMP_TARGET_KIND_OACC_KERNELS: + case GF_OMP_TARGET_KIND_OACC_SERIAL: case GF_OMP_TARGET_KIND_OACC_UPDATE: case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: case GF_OMP_TARGET_KIND_OACC_DECLARE: + case GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_PARALLELIZED: + case GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_GANG_SINGLE: data_region = false; break; case GF_OMP_TARGET_KIND_DATA: case GF_OMP_TARGET_KIND_OACC_DATA: case GF_OMP_TARGET_KIND_OACC_HOST_DATA: + case GF_OMP_TARGET_KIND_OACC_DATA_KERNELS: data_region = true; break; default: @@ -11091,7 +11833,10 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) case GOMP_MAP_FIRSTPRIVATE_REFERENCE: case GOMP_MAP_STRUCT: case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH: + case GOMP_MAP_DETACH: break; + case GOMP_MAP_IF_PRESENT: case GOMP_MAP_FORCE_ALLOC: case GOMP_MAP_FORCE_TO: case GOMP_MAP_FORCE_FROM: @@ -11100,6 +11845,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) case GOMP_MAP_FORCE_DEVICEPTR: case GOMP_MAP_DEVICE_RESIDENT: case GOMP_MAP_LINK: + case GOMP_MAP_FORCE_DETACH: gcc_assert (is_gimple_omp_oacc (stmt)); break; default: @@ -11153,10 +11899,20 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) continue; } + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH) + && is_omp_target (stmt)) + { + gcc_assert (maybe_lookup_field (c, ctx)); + map_cnt++; + continue; + } + if (!maybe_lookup_field (var, ctx)) continue; - /* Don't remap oacc parallel reduction variables, because the + /* Don't remap compute constructs' reduction variables, because the intermediate result must be local to each gang. */ if (offloaded && !(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_IN_REDUCTION (c))) @@ -11173,7 +11929,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) { gcc_assert (is_gimple_omp_oacc (ctx->stmt)); if (omp_is_reference (new_var) - && TREE_CODE (TREE_TYPE (new_var)) != POINTER_TYPE) + && (TREE_CODE (TREE_TYPE (new_var)) != POINTER_TYPE + || DECL_BY_REFERENCE (var))) { /* Create a local object to hold the instance value. */ @@ -11197,8 +11954,16 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) break; case OMP_CLAUSE_FIRSTPRIVATE: - if (is_oacc_parallel (ctx)) - goto oacc_firstprivate; + gcc_checking_assert (offloaded); + if (is_gimple_omp_oacc (ctx->stmt)) + { + /* No 'firstprivate' clauses on OpenACC 'kernels'. */ + gcc_checking_assert (!is_oacc_kernels (ctx)); + /* Likewise, on OpenACC 'kernels' decomposed parts. */ + gcc_checking_assert (!is_oacc_kernels_decomposed_part (ctx)); + + goto oacc_firstprivate; + } map_cnt++; var = OMP_CLAUSE_DECL (c); if (!omp_is_reference (var) @@ -11223,8 +11988,16 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) break; case OMP_CLAUSE_PRIVATE: + gcc_checking_assert (offloaded); if (is_gimple_omp_oacc (ctx->stmt)) - break; + { + /* No 'private' clauses on OpenACC 'kernels'. */ + gcc_checking_assert (!is_oacc_kernels (ctx)); + /* Likewise, on OpenACC 'kernels' decomposed parts. */ + gcc_checking_assert (!is_oacc_kernels_decomposed_part (ctx)); + + break; + } var = OMP_CLAUSE_DECL (c); if (is_variable_sized (var)) { @@ -11242,6 +12015,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) break; case OMP_CLAUSE_USE_DEVICE_PTR: + case OMP_CLAUSE_USE_DEVICE_ADDR: case OMP_CLAUSE_IS_DEVICE_PTR: var = OMP_CLAUSE_DECL (c); map_cnt++; @@ -11258,7 +12032,11 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) SET_DECL_VALUE_EXPR (new_var, x); DECL_HAS_VALUE_EXPR_P (new_var) = 1; } - else if (TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE) + else if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR + && !omp_is_reference (var) + && !omp_is_allocatable_or_ptr (var) + && !lang_hooks.decls.omp_array_data (var, true)) + || TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE) { tree new_var = lookup_decl (var, ctx); tree type = build_pointer_type (TREE_TYPE (var)); @@ -11379,14 +12157,28 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) gcc_assert (DECL_P (ovar2)); ovar = ovar2; } - if (!maybe_lookup_field (ovar, ctx)) + if (!maybe_lookup_field (ovar, ctx) + && !(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))) continue; } talign = TYPE_ALIGN_UNIT (TREE_TYPE (ovar)); if (DECL_P (ovar) && DECL_ALIGN_UNIT (ovar) > talign) talign = DECL_ALIGN_UNIT (ovar); - if (nc) + + if (nc + && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH) + && is_omp_target (stmt)) + { + var = lookup_decl_in_outer_ctx (ovar, ctx); + x = build_sender_ref (c, ctx); + gimplify_assign (x, build_fold_addr_expr (var), &ilist); + } + else if (nc) { var = lookup_decl_in_outer_ctx (ovar, ctx); x = build_sender_ref (ovar, ctx); @@ -11452,7 +12244,18 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) } else { - var = build_fold_addr_expr (var); + /* While MAP is handled explicitly by the FE, + for 'target update', only the identified is passed. */ + if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM + || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO) + && (omp_is_allocatable_or_ptr (var) + && omp_check_optional_argument (var, false))) + var = build_fold_indirect_ref (var); + else if ((OMP_CLAUSE_CODE (c) != OMP_CLAUSE_FROM + && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_TO) + || (!omp_is_allocatable_or_ptr (var) + && !omp_check_optional_argument (var, false))) + var = build_fold_addr_expr (var); gimplify_assign (x, var, &ilist); } } @@ -11461,7 +12264,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) { gcc_checking_assert (is_gimple_omp_oacc (ctx->stmt)); s = TREE_TYPE (ovar); - if (TREE_CODE (s) == REFERENCE_TYPE) + if (TREE_CODE (s) == REFERENCE_TYPE + || omp_check_optional_argument (ovar, false)) s = TREE_TYPE (s); s = TYPE_SIZE_UNIT (s); } @@ -11485,6 +12289,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) switch (tkind) { case GOMP_MAP_ALLOC: + case GOMP_MAP_IF_PRESENT: case GOMP_MAP_TO: case GOMP_MAP_FROM: case GOMP_MAP_TOFROM: @@ -11555,7 +12360,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) break; case OMP_CLAUSE_FIRSTPRIVATE: - if (is_oacc_parallel (ctx)) + if (is_gimple_omp_oacc (ctx->stmt)) goto oacc_firstprivate_map; ovar = OMP_CLAUSE_DECL (c); if (omp_is_reference (ovar)) @@ -11623,28 +12428,85 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) break; case OMP_CLAUSE_USE_DEVICE_PTR: + case OMP_CLAUSE_USE_DEVICE_ADDR: case OMP_CLAUSE_IS_DEVICE_PTR: ovar = OMP_CLAUSE_DECL (c); var = lookup_decl_in_outer_ctx (ovar, ctx); - x = build_sender_ref (ovar, ctx); - if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR) - tkind = GOMP_MAP_USE_DEVICE_PTR; + + if (lang_hooks.decls.omp_array_data (ovar, true)) + { + tkind = (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR + ? GOMP_MAP_USE_DEVICE_PTR : GOMP_MAP_FIRSTPRIVATE_INT); + x = build_sender_ref ((splay_tree_key) &DECL_NAME (ovar), ctx); + } + else if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR) + { + tkind = GOMP_MAP_USE_DEVICE_PTR; + x = build_sender_ref ((splay_tree_key) &DECL_UID (ovar), ctx); + } else - tkind = GOMP_MAP_FIRSTPRIVATE_INT; + { + tkind = GOMP_MAP_FIRSTPRIVATE_INT; + x = build_sender_ref (ovar, ctx); + } + + if (is_gimple_omp_oacc (ctx->stmt)) + { + gcc_assert (tkind == GOMP_MAP_USE_DEVICE_PTR); + + if (OMP_CLAUSE_USE_DEVICE_PTR_IF_PRESENT (c)) + tkind = GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT; + } + type = TREE_TYPE (ovar); - if (TREE_CODE (type) == ARRAY_TYPE) + if (lang_hooks.decls.omp_array_data (ovar, true)) + var = lang_hooks.decls.omp_array_data (ovar, false); + else if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR + && !omp_is_reference (ovar) + && !omp_is_allocatable_or_ptr (ovar)) + || TREE_CODE (type) == ARRAY_TYPE) var = build_fold_addr_expr (var); else { - if (omp_is_reference (ovar)) + if (omp_is_reference (ovar) + || omp_check_optional_argument (ovar, false) + || omp_is_allocatable_or_ptr (ovar)) { type = TREE_TYPE (type); - if (TREE_CODE (type) != ARRAY_TYPE) + if (TREE_CODE (type) != ARRAY_TYPE + && ((OMP_CLAUSE_CODE (c) != OMP_CLAUSE_USE_DEVICE_ADDR + && !omp_is_allocatable_or_ptr (ovar)) + || (omp_is_reference (ovar) + && omp_is_allocatable_or_ptr (ovar)))) var = build_simple_mem_ref (var); var = fold_convert (TREE_TYPE (x), var); } } - gimplify_assign (x, var, &ilist); + tree present; + present = omp_check_optional_argument (ovar, true); + if (present) + { + tree null_label = create_artificial_label (UNKNOWN_LOCATION); + tree notnull_label = create_artificial_label (UNKNOWN_LOCATION); + tree opt_arg_label = create_artificial_label (UNKNOWN_LOCATION); + tree new_x = unshare_expr (x); + gimplify_expr (&present, &ilist, NULL, is_gimple_val, + fb_rvalue); + gcond *cond = gimple_build_cond_from_tree (present, + notnull_label, + null_label); + gimple_seq_add_stmt (&ilist, cond); + gimple_seq_add_stmt (&ilist, gimple_build_label (null_label)); + gimplify_assign (new_x, null_pointer_node, &ilist); + gimple_seq_add_stmt (&ilist, gimple_build_goto (opt_arg_label)); + gimple_seq_add_stmt (&ilist, + gimple_build_label (notnull_label)); + gimplify_assign (x, var, &ilist); + gimple_seq_add_stmt (&ilist, + gimple_build_label (opt_arg_label)); + } + else + gimplify_assign (x, var, &ilist); s = size_int (0); purpose = size_int (map_idx++); CONSTRUCTOR_APPEND_ELT (vsize, purpose, s); @@ -11672,16 +12534,13 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) &initlist, true, NULL_TREE); gimple_seq_add_seq (&ilist, initlist); - tree clobber = build_constructor (TREE_TYPE (TREE_VEC_ELT (t, i)), - NULL); - TREE_THIS_VOLATILE (clobber) = 1; + tree clobber = build_clobber (TREE_TYPE (TREE_VEC_ELT (t, i))); gimple_seq_add_stmt (&olist, gimple_build_assign (TREE_VEC_ELT (t, i), clobber)); } - tree clobber = build_constructor (ctx->record_type, NULL); - TREE_THIS_VOLATILE (clobber) = 1; + tree clobber = build_clobber (ctx->record_type); gimple_seq_add_stmt (&olist, gimple_build_assign (ctx->sender_decl, clobber)); } @@ -11794,60 +12653,143 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) } break; case OMP_CLAUSE_USE_DEVICE_PTR: + case OMP_CLAUSE_USE_DEVICE_ADDR: case OMP_CLAUSE_IS_DEVICE_PTR: + tree new_var; + gimple_seq assign_body; + bool is_array_data; + bool do_optional_check; + assign_body = NULL; + do_optional_check = false; var = OMP_CLAUSE_DECL (c); - if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR) - x = build_sender_ref (var, ctx); + is_array_data = lang_hooks.decls.omp_array_data (var, true) != NULL; + + if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR) + x = build_sender_ref (is_array_data + ? (splay_tree_key) &DECL_NAME (var) + : (splay_tree_key) &DECL_UID (var), ctx); else x = build_receiver_ref (var, false, ctx); - if (is_variable_sized (var)) + + if (is_array_data) + { + bool is_ref = omp_is_reference (var); + do_optional_check = true; + /* First, we copy the descriptor data from the host; then + we update its data to point to the target address. */ + new_var = lookup_decl (var, ctx); + new_var = DECL_VALUE_EXPR (new_var); + tree v = new_var; + + if (is_ref) + { + var = build_fold_indirect_ref (var); + gimplify_expr (&var, &assign_body, NULL, is_gimple_val, + fb_rvalue); + v = create_tmp_var_raw (TREE_TYPE (var), get_name (var)); + gimple_add_tmp_var (v); + TREE_ADDRESSABLE (v) = 1; + gimple_seq_add_stmt (&assign_body, + gimple_build_assign (v, var)); + tree rhs = build_fold_addr_expr (v); + gimple_seq_add_stmt (&assign_body, + gimple_build_assign (new_var, rhs)); + } + else + gimple_seq_add_stmt (&assign_body, + gimple_build_assign (new_var, var)); + + tree v2 = lang_hooks.decls.omp_array_data (unshare_expr (v), false); + gcc_assert (v2); + gimplify_expr (&x, &assign_body, NULL, is_gimple_val, fb_rvalue); + gimple_seq_add_stmt (&assign_body, + gimple_build_assign (v2, x)); + } + else if (is_variable_sized (var)) { tree pvar = DECL_VALUE_EXPR (var); gcc_assert (TREE_CODE (pvar) == INDIRECT_REF); pvar = TREE_OPERAND (pvar, 0); gcc_assert (DECL_P (pvar)); - tree new_var = lookup_decl (pvar, ctx); - gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue); - gimple_seq_add_stmt (&new_body, + new_var = lookup_decl (pvar, ctx); + gimplify_expr (&x, &assign_body, NULL, is_gimple_val, fb_rvalue); + gimple_seq_add_stmt (&assign_body, gimple_build_assign (new_var, x)); } - else if (TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE) + else if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR + && !omp_is_reference (var) + && !omp_is_allocatable_or_ptr (var)) + || TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE) { - tree new_var = lookup_decl (var, ctx); + new_var = lookup_decl (var, ctx); new_var = DECL_VALUE_EXPR (new_var); gcc_assert (TREE_CODE (new_var) == MEM_REF); new_var = TREE_OPERAND (new_var, 0); gcc_assert (DECL_P (new_var)); - gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue); - gimple_seq_add_stmt (&new_body, + gimplify_expr (&x, &assign_body, NULL, is_gimple_val, fb_rvalue); + gimple_seq_add_stmt (&assign_body, gimple_build_assign (new_var, x)); } else { tree type = TREE_TYPE (var); - tree new_var = lookup_decl (var, ctx); + new_var = lookup_decl (var, ctx); if (omp_is_reference (var)) { type = TREE_TYPE (type); - if (TREE_CODE (type) != ARRAY_TYPE) + if (TREE_CODE (type) != ARRAY_TYPE + && (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_USE_DEVICE_ADDR + || (omp_is_reference (var) + && omp_is_allocatable_or_ptr (var)))) { tree v = create_tmp_var_raw (type, get_name (var)); gimple_add_tmp_var (v); TREE_ADDRESSABLE (v) = 1; x = fold_convert (type, x); - gimplify_expr (&x, &new_body, NULL, is_gimple_val, + gimplify_expr (&x, &assign_body, NULL, is_gimple_val, fb_rvalue); - gimple_seq_add_stmt (&new_body, + gimple_seq_add_stmt (&assign_body, gimple_build_assign (v, x)); x = build_fold_addr_expr (v); + do_optional_check = true; } } new_var = DECL_VALUE_EXPR (new_var); x = fold_convert (TREE_TYPE (new_var), x); - gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue); - gimple_seq_add_stmt (&new_body, + gimplify_expr (&x, &assign_body, NULL, is_gimple_val, fb_rvalue); + gimple_seq_add_stmt (&assign_body, gimple_build_assign (new_var, x)); } + tree present; + present = (do_optional_check + ? omp_check_optional_argument (OMP_CLAUSE_DECL (c), true) + : NULL_TREE); + if (present) + { + tree null_label = create_artificial_label (UNKNOWN_LOCATION); + tree notnull_label = create_artificial_label (UNKNOWN_LOCATION); + tree opt_arg_label = create_artificial_label (UNKNOWN_LOCATION); + glabel *null_glabel = gimple_build_label (null_label); + glabel *notnull_glabel = gimple_build_label (notnull_label); + ggoto *opt_arg_ggoto = gimple_build_goto (opt_arg_label); + gimplify_expr (&x, &new_body, NULL, is_gimple_val, + fb_rvalue); + gimplify_expr (&present, &new_body, NULL, is_gimple_val, + fb_rvalue); + gcond *cond = gimple_build_cond_from_tree (present, + notnull_label, + null_label); + gimple_seq_add_stmt (&new_body, cond); + gimple_seq_add_stmt (&new_body, null_glabel); + gimplify_assign (new_var, null_pointer_node, &new_body); + gimple_seq_add_stmt (&new_body, opt_arg_ggoto); + gimple_seq_add_stmt (&new_body, notnull_glabel); + gimple_seq_add_seq (&new_body, assign_body); + gimple_seq_add_stmt (&new_body, + gimple_build_label (opt_arg_label)); + } + else + gimple_seq_add_seq (&new_body, assign_body); break; } /* Handle GOMP_MAP_FIRSTPRIVATE_{POINTER,REFERENCE} in second pass, @@ -12022,7 +12964,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) gimple_seq fork_seq = NULL; gimple_seq join_seq = NULL; - if (is_oacc_parallel (ctx)) + if (offloaded && is_gimple_omp_oacc (ctx->stmt)) { /* If there are reductions on the offloaded region itself, treat them as a dummy GANG loop. */ @@ -12104,22 +13046,19 @@ lower_omp_teams (gimple_stmt_iterator *gsi_p, omp_context *ctx) lower_omp (gimple_omp_body_ptr (teams_stmt), ctx); lower_reduction_clauses (gimple_omp_teams_clauses (teams_stmt), &olist, NULL, ctx); - if (!gimple_omp_teams_grid_phony (teams_stmt)) - { - gimple_seq_add_stmt (&bind_body, teams_stmt); - location_t loc = gimple_location (teams_stmt); - tree decl = builtin_decl_explicit (BUILT_IN_GOMP_TEAMS); - gimple *call = gimple_build_call (decl, 2, num_teams, thread_limit); - gimple_set_location (call, loc); - gimple_seq_add_stmt (&bind_body, call); - } + gimple_seq_add_stmt (&bind_body, teams_stmt); + + location_t loc = gimple_location (teams_stmt); + tree decl = builtin_decl_explicit (BUILT_IN_GOMP_TEAMS); + gimple *call = gimple_build_call (decl, 2, num_teams, thread_limit); + gimple_set_location (call, loc); + gimple_seq_add_stmt (&bind_body, call); gimple_seq_add_seq (&bind_body, gimple_omp_body (teams_stmt)); gimple_omp_set_body (teams_stmt, NULL); gimple_seq_add_seq (&bind_body, olist); gimple_seq_add_seq (&bind_body, dlist); - if (!gimple_omp_teams_grid_phony (teams_stmt)) - gimple_seq_add_stmt (&bind_body, gimple_build_omp_return (true)); + gimple_seq_add_stmt (&bind_body, gimple_build_omp_return (true)); gimple_bind_set_body (bind, bind_body); pop_gimplify_context (bind); @@ -12130,18 +13069,6 @@ lower_omp_teams (gimple_stmt_iterator *gsi_p, omp_context *ctx) TREE_USED (block) = 1; } -/* Expand code within an artificial GIMPLE_OMP_GRID_BODY OMP construct. */ - -static void -lower_omp_grid_body (gimple_stmt_iterator *gsi_p, omp_context *ctx) -{ - gimple *stmt = gsi_stmt (*gsi_p); - lower_omp (gimple_omp_body_ptr (stmt), ctx); - gimple_seq_add_stmt (gimple_omp_body_ptr (stmt), - gimple_build_omp_return (false)); -} - - /* Callback for lower_omp_1. Return non-NULL if *tp needs to be regimplified. If DATA is non-NULL, lower_omp_1 is outside of OMP context, but with task_shared_vars set. */ @@ -12360,11 +13287,6 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context *ctx) else lower_omp_teams (gsi_p, ctx); break; - case GIMPLE_OMP_GRID_BODY: - ctx = maybe_lookup_ctx (stmt); - gcc_assert (ctx); - lower_omp_grid_body (gsi_p, ctx); - break; case GIMPLE_CALL: tree fndecl; call_stmt = as_a (stmt); @@ -12522,9 +13444,6 @@ execute_lower_omp (void) body = gimple_body (current_function_decl); - if (hsa_gen_requested_p ()) - omp_grid_gridify_all_targets (&body); - scan_omp (&body, NULL); gcc_assert (taskreg_nesting_level == 0); FOR_EACH_VEC_ELT (taskreg_contexts, i, ctx) @@ -12546,6 +13465,7 @@ execute_lower_omp (void) all_contexts = NULL; } BITMAP_FREE (task_shared_vars); + BITMAP_FREE (global_nonaddressable_vars); /* If current function is a method, remove artificial dummy VAR_DECL created for non-static data member privatization, they aren't needed for