diff options
Diffstat (limited to 'gcc/c-family/c-omp.cc')
| -rw-r--r-- | gcc/c-family/c-omp.cc | 865 |
1 files changed, 803 insertions, 62 deletions
diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc index a92c6e3..a3cd1ae 100644 --- a/gcc/c-family/c-omp.cc +++ b/gcc/c-family/c-omp.cc @@ -764,9 +764,7 @@ c_finish_omp_depobj (location_t loc, tree depobj, kind = OMP_CLAUSE_DEPEND_KIND (clause); t = OMP_CLAUSE_DECL (clause); gcc_assert (t); - if (TREE_CODE (t) == TREE_LIST - && TREE_PURPOSE (t) - && TREE_CODE (TREE_PURPOSE (t)) == TREE_VEC) + if (OMP_ITERATOR_DECL_P (t)) { error_at (OMP_CLAUSE_LOCATION (clause), "%<iterator%> modifier may not be specified on " @@ -2178,6 +2176,7 @@ c_omp_split_clauses (location_t loc, enum tree_code code, case OMP_CLAUSE_HAS_DEVICE_ADDR: case OMP_CLAUSE_DEFAULTMAP: case OMP_CLAUSE_DEPEND: + case OMP_CLAUSE_USES_ALLOCATORS: s = C_OMP_CLAUSE_SPLIT_TARGET; break; case OMP_CLAUSE_DOACROSS: @@ -3546,7 +3545,9 @@ c_omp_address_inspector::map_supported_p () || TREE_CODE (t) == POINTER_PLUS_EXPR || TREE_CODE (t) == NON_LVALUE_EXPR || TREE_CODE (t) == OMP_ARRAY_SECTION - || TREE_CODE (t) == NOP_EXPR) + || TREE_CODE (t) == NOP_EXPR + || TREE_CODE (t) == VIEW_CONVERT_EXPR + || TREE_CODE (t) == ADDR_EXPR) if (TREE_CODE (t) == COMPOUND_EXPR) t = TREE_OPERAND (t, 1); else @@ -3641,11 +3642,12 @@ c_omp_address_inspector::maybe_zero_length_array_section (tree clause) expression types here, because e.g. you can't have an array of references. */ -static tree -omp_expand_access_chain (tree c, tree expr, vec<omp_addr_token *> &addr_tokens, - unsigned *idx, c_omp_region_type ort) +static tree * +omp_expand_access_chain (tree *pc, tree expr, + vec<omp_addr_token *> &addr_tokens, unsigned *idx, c_omp_region_type ort) { using namespace omp_addr_tokenizer; + tree c = *pc; location_t loc = OMP_CLAUSE_LOCATION (c); unsigned i = *idx; tree c2 = NULL_TREE; @@ -3688,39 +3690,114 @@ omp_expand_access_chain (tree c, tree expr, vec<omp_addr_token *> &addr_tokens, break; default: - return error_mark_node; + return NULL; } if (c2) { OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c); OMP_CLAUSE_CHAIN (c) = c2; - c = c2; + pc = &OMP_CLAUSE_CHAIN (c); } *idx = ++i; if (i < addr_tokens.length () && addr_tokens[i]->type == ACCESS_METHOD) - return omp_expand_access_chain (c, expr, addr_tokens, idx, ort); + return omp_expand_access_chain (pc, expr, addr_tokens, idx, ort); - return c; + return pc; +} + +static tree * +omp_expand_grid_dim (location_t loc, tree *pc, tree decl) +{ + if (TREE_CODE (decl) == OMP_ARRAY_SECTION) + pc = omp_expand_grid_dim (loc, pc, TREE_OPERAND (decl, 0)); + else + return pc; + + tree c = *pc; + tree low_bound = TREE_OPERAND (decl, 1); + tree length = TREE_OPERAND (decl, 2); + tree stride = TREE_OPERAND (decl, 3); + + tree cd = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (cd, GOMP_MAP_GRID_DIM); + OMP_CLAUSE_DECL (cd) = unshare_expr (low_bound); + OMP_CLAUSE_SIZE (cd) = unshare_expr (length); + + if (stride && !integer_onep (stride)) + { + tree cs = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (cs, GOMP_MAP_GRID_STRIDE); + OMP_CLAUSE_DECL (cs) = unshare_expr (stride); + + OMP_CLAUSE_CHAIN (cs) = OMP_CLAUSE_CHAIN (c); + OMP_CLAUSE_CHAIN (cd) = cs; + OMP_CLAUSE_CHAIN (c) = cd; + pc = &OMP_CLAUSE_CHAIN (cd); + } + else + { + OMP_CLAUSE_CHAIN (cd) = OMP_CLAUSE_CHAIN (c); + OMP_CLAUSE_CHAIN (c) = cd; + pc = &OMP_CLAUSE_CHAIN (c); + } + + return pc; +} + +tree * +omp_handle_noncontig_array (location_t loc, tree *pc, tree c, tree base) +{ + tree type; + + if (POINTER_TYPE_P (TREE_TYPE (base))) + type = TREE_TYPE (TREE_TYPE (base)); + else + type = strip_array_types (TREE_TYPE (base)); + + tree c_map = build_omp_clause (loc, OMP_CLAUSE_MAP); + + OMP_CLAUSE_DECL (c_map) = unshare_expr (base); + /* Use the element size (or pointed-to type size) here. */ + OMP_CLAUSE_SIZE (c_map) = TYPE_SIZE_UNIT (type); + + switch (OMP_CLAUSE_CODE (c)) + { + case OMP_CLAUSE_TO: + OMP_CLAUSE_SET_MAP_KIND (c_map, GOMP_MAP_TO_GRID); + break; + case OMP_CLAUSE_FROM: + OMP_CLAUSE_SET_MAP_KIND (c_map, GOMP_MAP_FROM_GRID); + break; + default: + gcc_unreachable (); + } + + OMP_CLAUSE_CHAIN (c_map) = OMP_CLAUSE_CHAIN (c); + + *pc = c_map; + + return omp_expand_grid_dim (loc, pc, OMP_CLAUSE_DECL (c)); } /* Translate "array_base_decl access_method" to OMP mapping clauses. */ -tree -c_omp_address_inspector::expand_array_base (tree c, +tree * +c_omp_address_inspector::expand_array_base (tree *pc, vec<omp_addr_token *> &addr_tokens, tree expr, unsigned *idx, c_omp_region_type ort) { using namespace omp_addr_tokenizer; + tree c = *pc; location_t loc = OMP_CLAUSE_LOCATION (c); int i = *idx; tree decl = addr_tokens[i + 1]->expr; bool decl_p = DECL_P (decl); - bool declare_target_p = (decl_p + bool declare_target_p = (DECL_P (decl) && is_global_var (decl) && lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl))); @@ -3731,6 +3808,7 @@ c_omp_address_inspector::expand_array_base (tree c, unsigned consume_tokens = 2; bool target_p = (ort & C_ORT_TARGET) != 0; bool openmp_p = (ort & C_ORT_OMP) != 0; + unsigned acc = i + 1; gcc_assert (i == 0); @@ -3741,10 +3819,18 @@ c_omp_address_inspector::expand_array_base (tree c, { i += 2; *idx = i; - return c; + return pc; + } + + if (!map_p && chain_p) + { + /* See comment in c_omp_address_inspector::expand_component_selector. */ + while (acc + 1 < addr_tokens.length () + && addr_tokens[acc + 1]->type == ACCESS_METHOD) + acc++; } - switch (addr_tokens[i + 1]->u.access_kind) + switch (addr_tokens[acc]->u.access_kind) { case ACCESS_DIRECT: if (decl_p && !target_p) @@ -3867,7 +3953,8 @@ c_omp_address_inspector::expand_array_base (tree c, /* The code handling "firstprivatize_array_bases" in gimplify.cc is relevant here. What do we need to create for arrays at this stage? (This condition doesn't feel quite right. FIXME?) */ - if (!target_p + if (openmp_p + && !target_p && (TREE_CODE (TREE_TYPE (addr_tokens[i + 1]->expr)) == ARRAY_TYPE)) break; @@ -3878,7 +3965,7 @@ c_omp_address_inspector::expand_array_base (tree c, virtual_origin); tree data_addr = omp_accessed_addr (addr_tokens, i + 1, expr); c2 = build_omp_clause (loc, OMP_CLAUSE_MAP); - if (decl_p && target_p) + if (decl_p && (!openmp_p || target_p)) { /* See comment for ACCESS_INDEXED_REF_TO_ARRAY above. */ enum gomp_map_kind k = chain_p ? GOMP_MAP_POINTER @@ -3934,9 +4021,11 @@ c_omp_address_inspector::expand_array_base (tree c, tree data_addr = omp_accessed_addr (addr_tokens, last_access, expr); c2 = build_omp_clause (loc, OMP_CLAUSE_MAP); /* For OpenACC, use FIRSTPRIVATE_POINTER for decls even on non-compute - regions (e.g. "acc data" constructs). It'll be removed anyway in - gimplify.cc, but doing it this way maintains diagnostic - behaviour. */ + regions (e.g. "acc data" constructs). It is used during "lexical + inheritance" of mapping clauses on enclosed target + (parallel/serial/kernels) regions, i.e. creating "present" mappings + for sections of pointer-based arrays. It's also used for + diagnostics. */ if (decl_p && (target_p || !openmp_p) && !chain_p && !declare_target_p) OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER); else @@ -4013,9 +4102,43 @@ c_omp_address_inspector::expand_array_base (tree c, } break; + case ACCESS_NONCONTIG_ARRAY: + { + gcc_assert (!map_p); + + tree base = addr_tokens[acc]->expr; + + if (decl_p) + c_common_mark_addressable_vec (base); + + pc = omp_handle_noncontig_array (loc, pc, c, base); + consume_tokens = (acc + 1) - i; + chain_p = false; + } + break; + + case ACCESS_NONCONTIG_REF_TO_ARRAY: + { + gcc_assert (!map_p); + + if (decl_p) + c_common_mark_addressable_vec (addr_tokens[acc]->expr); + + /* Or here. */ + gcc_assert (!chain_p); + + tree base = addr_tokens[i + 1]->expr; + base = convert_from_reference (base); + + pc = omp_handle_noncontig_array (loc, pc, c, base); + consume_tokens = (acc + 1) - i; + chain_p = false; + } + break; + default: *idx = i + consume_tokens; - return error_mark_node; + return NULL; } if (c3) @@ -4028,43 +4151,65 @@ c_omp_address_inspector::expand_array_base (tree c, OMP_CLAUSE_MAP_IMPLICIT (c2) = 1; OMP_CLAUSE_MAP_IMPLICIT (c3) = 1; } - c = c3; + pc = &OMP_CLAUSE_CHAIN (c2); } else if (c2) { + if (OMP_CLAUSE_MAP_READONLY (c)) + OMP_CLAUSE_MAP_POINTS_TO_READONLY (c2) = 1; OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c); OMP_CLAUSE_CHAIN (c) = c2; if (implicit_p) OMP_CLAUSE_MAP_IMPLICIT (c2) = 1; - c = c2; + pc = &OMP_CLAUSE_CHAIN (c); } i += consume_tokens; *idx = i; if (chain_p && map_p) - return omp_expand_access_chain (c, expr, addr_tokens, idx, ort); + return omp_expand_access_chain (pc, expr, addr_tokens, idx, ort); - return c; + return pc; } /* Translate "component_selector access_method" to OMP mapping clauses. */ -tree -c_omp_address_inspector::expand_component_selector (tree c, +tree * +c_omp_address_inspector::expand_component_selector (tree *pc, vec<omp_addr_token *> &addr_tokens, tree expr, unsigned *idx, c_omp_region_type ort) { using namespace omp_addr_tokenizer; + tree c = *pc; location_t loc = OMP_CLAUSE_LOCATION (c); unsigned i = *idx; tree c2 = NULL_TREE, c3 = NULL_TREE; bool chain_p = omp_access_chain_p (addr_tokens, i + 1); bool map_p = OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP; + unsigned acc = i + 1; + + if (!map_p && chain_p) + { + /* We have a non-map clause (i.e. to/from for an "update" directive), + and we might have a noncontiguous array section at the end of a + chain of other accesses, e.g. pointer indirections like this: + + struct_base_decl access_pointer access_pointer component_selector + access_pointer access_pointer access_noncontig_array - switch (addr_tokens[i + 1]->u.access_kind) + We only need to process the last access in this case, so skip + over previous accesses. */ + + while (acc + 1 < addr_tokens.length () + && addr_tokens[acc + 1]->type == ACCESS_METHOD) + acc++; + chain_p = false; + } + + switch (addr_tokens[acc]->u.access_kind) { case ACCESS_DIRECT: case ACCESS_INDEXED_ARRAY: @@ -4074,7 +4219,7 @@ c_omp_address_inspector::expand_component_selector (tree c, { /* Copy the referenced object. Note that we also do this for !MAP_P clauses. */ - tree obj = convert_from_reference (addr_tokens[i + 1]->expr); + tree obj = convert_from_reference (addr_tokens[acc]->expr); OMP_CLAUSE_DECL (c) = obj; OMP_CLAUSE_SIZE (c) = TYPE_SIZE_UNIT (TREE_TYPE (obj)); @@ -4083,7 +4228,7 @@ c_omp_address_inspector::expand_component_selector (tree c, c2 = build_omp_clause (loc, OMP_CLAUSE_MAP); OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH); - OMP_CLAUSE_DECL (c2) = addr_tokens[i + 1]->expr; + OMP_CLAUSE_DECL (c2) = addr_tokens[acc]->expr; OMP_CLAUSE_SIZE (c2) = size_zero_node; } break; @@ -4094,15 +4239,15 @@ c_omp_address_inspector::expand_component_selector (tree c, break; tree virtual_origin - = convert_from_reference (addr_tokens[i + 1]->expr); + = convert_from_reference (addr_tokens[acc]->expr); virtual_origin = build_fold_addr_expr (virtual_origin); virtual_origin = fold_convert_loc (loc, ptrdiff_type_node, virtual_origin); - tree data_addr = omp_accessed_addr (addr_tokens, i + 1, expr); + tree data_addr = omp_accessed_addr (addr_tokens, acc, expr); c2 = build_omp_clause (loc, OMP_CLAUSE_MAP); OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH); - OMP_CLAUSE_DECL (c2) = addr_tokens[i + 1]->expr; + OMP_CLAUSE_DECL (c2) = addr_tokens[acc]->expr; OMP_CLAUSE_SIZE (c2) = fold_build2_loc (loc, MINUS_EXPR, ptrdiff_type_node, fold_convert_loc (loc, ptrdiff_type_node, @@ -4119,12 +4264,12 @@ c_omp_address_inspector::expand_component_selector (tree c, tree virtual_origin = fold_convert_loc (loc, ptrdiff_type_node, - addr_tokens[i + 1]->expr); - tree data_addr = omp_accessed_addr (addr_tokens, i + 1, expr); + addr_tokens[acc]->expr); + tree data_addr = omp_accessed_addr (addr_tokens, acc, expr); c2 = build_omp_clause (loc, OMP_CLAUSE_MAP); OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH); - OMP_CLAUSE_DECL (c2) = addr_tokens[i + 1]->expr; + OMP_CLAUSE_DECL (c2) = addr_tokens[acc]->expr; OMP_CLAUSE_SIZE (c2) = fold_build2_loc (loc, MINUS_EXPR, ptrdiff_type_node, fold_convert_loc (loc, ptrdiff_type_node, @@ -4139,10 +4284,10 @@ c_omp_address_inspector::expand_component_selector (tree c, if (!map_p) break; - tree ptr = convert_from_reference (addr_tokens[i + 1]->expr); + tree ptr = convert_from_reference (addr_tokens[acc]->expr); tree virtual_origin = fold_convert_loc (loc, ptrdiff_type_node, ptr); - tree data_addr = omp_accessed_addr (addr_tokens, i + 1, expr); + tree data_addr = omp_accessed_addr (addr_tokens, acc, expr); /* Attach the pointer... */ c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); @@ -4157,14 +4302,39 @@ c_omp_address_inspector::expand_component_selector (tree c, /* ...and also the reference. */ c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); OMP_CLAUSE_SET_MAP_KIND (c3, GOMP_MAP_ATTACH_DETACH); - OMP_CLAUSE_DECL (c3) = addr_tokens[i + 1]->expr; + OMP_CLAUSE_DECL (c3) = addr_tokens[acc]->expr; OMP_CLAUSE_SIZE (c3) = size_zero_node; } break; + case ACCESS_NONCONTIG_ARRAY: + { + gcc_assert (!map_p); + + /* We don't expect to see further accesses here. */ + gcc_assert (!chain_p); + + pc = omp_handle_noncontig_array (loc, pc, c, addr_tokens[acc]->expr); + } + break; + + case ACCESS_NONCONTIG_REF_TO_ARRAY: + { + gcc_assert (!map_p); + + /* Or here. */ + gcc_assert (!chain_p); + + tree base = addr_tokens[acc]->expr; + base = convert_from_reference (base); + + pc = omp_handle_noncontig_array (loc, pc, c, base); + } + break; + default: - *idx = i + 2; - return error_mark_node; + *idx = acc + 1; + return NULL; } if (c3) @@ -4172,29 +4342,30 @@ c_omp_address_inspector::expand_component_selector (tree c, OMP_CLAUSE_CHAIN (c3) = OMP_CLAUSE_CHAIN (c); OMP_CLAUSE_CHAIN (c2) = c3; OMP_CLAUSE_CHAIN (c) = c2; - c = c3; + pc = &OMP_CLAUSE_CHAIN (c2); } else if (c2) { + if (OMP_CLAUSE_MAP_READONLY (c)) + OMP_CLAUSE_MAP_POINTS_TO_READONLY (c2) = 1; OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c); OMP_CLAUSE_CHAIN (c) = c2; - c = c2; + pc = &OMP_CLAUSE_CHAIN (c); } - i += 2; - *idx = i; + *idx = acc + 1; if (chain_p && map_p) - return omp_expand_access_chain (c, expr, addr_tokens, idx, ort); + return omp_expand_access_chain (pc, expr, addr_tokens, idx, ort); - return c; + return pc; } /* Expand a map clause into a group of mapping clauses, creating nodes to attach/detach pointers and so forth as necessary. */ -tree -c_omp_address_inspector::expand_map_clause (tree c, tree expr, +tree * +c_omp_address_inspector::expand_map_clause (tree *pc, tree expr, vec<omp_addr_token *> &addr_tokens, c_omp_region_type ort) { @@ -4210,18 +4381,18 @@ c_omp_address_inspector::expand_map_clause (tree c, tree expr, && addr_tokens[i]->u.structure_base_kind == BASE_DECL && addr_tokens[i + 1]->type == ACCESS_METHOD) { - c = expand_array_base (c, addr_tokens, expr, &i, ort); - if (c == error_mark_node) - return error_mark_node; + pc = expand_array_base (pc, addr_tokens, expr, &i, ort); + if (pc == NULL) + return NULL; } else if (remaining >= 2 && addr_tokens[i]->type == ARRAY_BASE && addr_tokens[i]->u.structure_base_kind == BASE_ARBITRARY_EXPR && addr_tokens[i + 1]->type == ACCESS_METHOD) { - c = expand_array_base (c, addr_tokens, expr, &i, ort); - if (c == error_mark_node) - return error_mark_node; + pc = expand_array_base (pc, addr_tokens, expr, &i, ort); + if (pc == NULL) + return NULL; } else if (remaining >= 2 && addr_tokens[i]->type == STRUCTURE_BASE @@ -4248,18 +4419,18 @@ c_omp_address_inspector::expand_map_clause (tree c, tree expr, i++; break; default: - return error_mark_node; + return NULL; } } else if (remaining >= 2 && addr_tokens[i]->type == COMPONENT_SELECTOR && addr_tokens[i + 1]->type == ACCESS_METHOD) { - c = expand_component_selector (c, addr_tokens, expr, &i, ort); + pc = expand_component_selector (pc, addr_tokens, expr, &i, ort); /* We used 'expr', so these must have been the last tokens. */ gcc_assert (i == length); - if (c == error_mark_node) - return error_mark_node; + if (pc == NULL) + return NULL; } else if (remaining >= 3 && addr_tokens[i]->type == COMPONENT_SELECTOR @@ -4277,9 +4448,579 @@ c_omp_address_inspector::expand_map_clause (tree c, tree expr, } if (i == length) - return c; + return pc; - return error_mark_node; + return NULL; +} + +/* Given a mapper function MAPPER_FN, recursively scan through the map clauses + for that mapper, and if any of those should use a (named or unnamed) mapper + themselves, add it to MLIST. */ + +void +c_omp_find_nested_mappers (omp_mapper_list<tree> *mlist, tree mapper_fn) +{ + tree mapper = lang_hooks.decls.omp_extract_mapper_directive (mapper_fn); + tree mapper_name = NULL_TREE; + + if (mapper == error_mark_node) + return; + + gcc_assert (TREE_CODE (mapper) == OMP_DECLARE_MAPPER); + + for (tree clause = OMP_DECLARE_MAPPER_CLAUSES (mapper); + clause; + clause = OMP_CLAUSE_CHAIN (clause)) + { + tree expr = OMP_CLAUSE_DECL (clause); + enum gomp_map_kind clause_kind = OMP_CLAUSE_MAP_KIND (clause); + tree elem_type; + + if (clause_kind == GOMP_MAP_PUSH_MAPPER_NAME) + { + mapper_name = expr; + continue; + } + else if (clause_kind == GOMP_MAP_POP_MAPPER_NAME) + { + mapper_name = NULL_TREE; + continue; + } + + gcc_assert (TREE_CODE (expr) != TREE_LIST); + if (TREE_CODE (expr) == OMP_ARRAY_SECTION) + { + while (TREE_CODE (expr) == OMP_ARRAY_SECTION) + expr = TREE_OPERAND (expr, 0); + + elem_type = TREE_TYPE (expr); + } + else + elem_type = TREE_TYPE (expr); + + /* This might be too much... or not enough? */ + while (TREE_CODE (elem_type) == ARRAY_TYPE + || TREE_CODE (elem_type) == POINTER_TYPE + || TREE_CODE (elem_type) == REFERENCE_TYPE) + elem_type = TREE_TYPE (elem_type); + + elem_type = TYPE_MAIN_VARIANT (elem_type); + + if (AGGREGATE_TYPE_P (elem_type) + && !mlist->contains (mapper_name, elem_type)) + { + tree nested_mapper_fn + = lang_hooks.decls.omp_mapper_lookup (mapper_name, elem_type); + + if (nested_mapper_fn) + { + mlist->add_mapper (mapper_name, elem_type, nested_mapper_fn); + c_omp_find_nested_mappers (mlist, nested_mapper_fn); + } + else if (mapper_name) + { + error ("mapper %qE not found for type %qT", mapper_name, + elem_type); + continue; + } + } + } +} + +struct remap_mapper_decl_info +{ + tree dummy_var; + tree expr; +}; + +/* Helper for rewriting DUMMY_VAR into EXPR in a map clause decl. */ + +static tree +remap_mapper_decl_1 (tree *tp, int *walk_subtrees, void *data) +{ + remap_mapper_decl_info *map_info = (remap_mapper_decl_info *) data; + + if (operand_equal_p (*tp, map_info->dummy_var)) + { + *tp = map_info->expr; + *walk_subtrees = 0; + } + + return NULL_TREE; +} + +static enum gomp_map_kind +omp_split_map_kind (enum gomp_map_kind op, bool *force_p, bool *always_p, + bool *present_p) +{ + *force_p = *always_p = *present_p = false; + + switch (op) + { + case GOMP_MAP_FORCE_ALLOC: + case GOMP_MAP_FORCE_TO: + case GOMP_MAP_FORCE_FROM: + case GOMP_MAP_FORCE_TOFROM: + case GOMP_MAP_FORCE_PRESENT: + *force_p = true; + break; + case GOMP_MAP_ALWAYS_TO: + case GOMP_MAP_ALWAYS_FROM: + case GOMP_MAP_ALWAYS_TOFROM: + *always_p = true; + break; + case GOMP_MAP_ALWAYS_PRESENT_TO: + case GOMP_MAP_ALWAYS_PRESENT_FROM: + case GOMP_MAP_ALWAYS_PRESENT_TOFROM: + *always_p = true; + /* Fallthrough. */ + case GOMP_MAP_PRESENT_ALLOC: + case GOMP_MAP_PRESENT_TO: + case GOMP_MAP_PRESENT_FROM: + case GOMP_MAP_PRESENT_TOFROM: + *present_p = true; + break; + default: + ; + } + + switch (op) + { + case GOMP_MAP_ALLOC: + case GOMP_MAP_FORCE_ALLOC: + case GOMP_MAP_PRESENT_ALLOC: + return GOMP_MAP_ALLOC; + case GOMP_MAP_TO: + case GOMP_MAP_FORCE_TO: + case GOMP_MAP_ALWAYS_TO: + case GOMP_MAP_PRESENT_TO: + case GOMP_MAP_ALWAYS_PRESENT_TO: + return GOMP_MAP_TO; + case GOMP_MAP_FROM: + case GOMP_MAP_FORCE_FROM: + case GOMP_MAP_ALWAYS_FROM: + case GOMP_MAP_PRESENT_FROM: + case GOMP_MAP_ALWAYS_PRESENT_FROM: + return GOMP_MAP_FROM; + case GOMP_MAP_TOFROM: + case GOMP_MAP_FORCE_TOFROM: + case GOMP_MAP_ALWAYS_TOFROM: + case GOMP_MAP_PRESENT_TOFROM: + case GOMP_MAP_ALWAYS_PRESENT_TOFROM: + return GOMP_MAP_TOFROM; + default: + ; + } + + return op; +} + +static enum gomp_map_kind +omp_join_map_kind (enum gomp_map_kind op, bool force_p, bool always_p, + bool present_p) +{ + gcc_assert (!force_p || !(always_p || present_p)); + + switch (op) + { + case GOMP_MAP_ALLOC: + if (force_p) + return GOMP_MAP_FORCE_ALLOC; + else if (present_p) + return GOMP_MAP_PRESENT_ALLOC; + break; + + case GOMP_MAP_TO: + if (force_p) + return GOMP_MAP_FORCE_TO; + else if (always_p && present_p) + return GOMP_MAP_ALWAYS_PRESENT_TO; + else if (always_p) + return GOMP_MAP_ALWAYS_TO; + else if (present_p) + return GOMP_MAP_PRESENT_TO; + break; + + case GOMP_MAP_FROM: + if (force_p) + return GOMP_MAP_FORCE_FROM; + else if (always_p && present_p) + return GOMP_MAP_ALWAYS_PRESENT_FROM; + else if (always_p) + return GOMP_MAP_ALWAYS_FROM; + else if (present_p) + return GOMP_MAP_PRESENT_FROM; + break; + + case GOMP_MAP_TOFROM: + if (force_p) + return GOMP_MAP_FORCE_TOFROM; + else if (always_p && present_p) + return GOMP_MAP_ALWAYS_PRESENT_TOFROM; + else if (always_p) + return GOMP_MAP_ALWAYS_TOFROM; + else if (present_p) + return GOMP_MAP_PRESENT_TOFROM; + break; + + default: + ; + } + + return op; +} + +/* Map kind decay (OpenMP 5.2, 5.8.8 "declare mapper Directive"). Return the + map kind to use given MAPPER_KIND specified in the mapper and INVOKED_AS + specified on the clause that invokes the mapper. See also + fortran/trans-openmp.cc:omp_map_decayed_kind. */ + +static enum gomp_map_kind +omp_map_decayed_kind (enum gomp_map_kind mapper_kind, + enum gomp_map_kind invoked_as, bool exit_p) +{ + if (invoked_as == GOMP_MAP_RELEASE || invoked_as == GOMP_MAP_DELETE) + return invoked_as; + + bool force_p, always_p, present_p; + + invoked_as = omp_split_map_kind (invoked_as, &force_p, &always_p, &present_p); + gomp_map_kind decay_to; + + switch (mapper_kind) + { + case GOMP_MAP_ALLOC: + if (exit_p && invoked_as == GOMP_MAP_FROM) + decay_to = GOMP_MAP_RELEASE; + else + decay_to = GOMP_MAP_ALLOC; + break; + + case GOMP_MAP_TO: + if (invoked_as == GOMP_MAP_FROM) + decay_to = exit_p ? GOMP_MAP_RELEASE : GOMP_MAP_ALLOC; + else if (invoked_as == GOMP_MAP_ALLOC) + decay_to = GOMP_MAP_ALLOC; + else + decay_to = GOMP_MAP_TO; + break; + + case GOMP_MAP_FROM: + if (invoked_as == GOMP_MAP_ALLOC || invoked_as == GOMP_MAP_TO) + decay_to = GOMP_MAP_ALLOC; + else + decay_to = GOMP_MAP_FROM; + break; + + case GOMP_MAP_TOFROM: + case GOMP_MAP_UNSET: + decay_to = invoked_as; + break; + + default: + gcc_unreachable (); + } + + return omp_join_map_kind (decay_to, force_p, always_p, present_p); +} + +/* Return a name to use for a "basic" map kind, e.g. as output from + omp_split_map_kind above. */ + +static const char * +omp_basic_map_kind_name (enum gomp_map_kind kind) +{ + switch (kind) + { + case GOMP_MAP_ALLOC: + return "alloc"; + case GOMP_MAP_TO: + return "to"; + case GOMP_MAP_FROM: + return "from"; + case GOMP_MAP_TOFROM: + return "tofrom"; + case GOMP_MAP_RELEASE: + return "release"; + case GOMP_MAP_DELETE: + return "delete"; + default: + gcc_unreachable (); + } +} + +/* Instantiate a mapper MAPPER for expression EXPR, adding new clauses to + OUTLIST. OUTER_KIND is the mapping kind to use if not already specified in + the mapper declaration. */ + +static tree * +omp_instantiate_mapper (location_t loc, tree *outlist, tree mapper, tree expr, + enum gomp_map_kind outer_kind, + enum c_omp_region_type ort) +{ + tree clauses = OMP_DECLARE_MAPPER_CLAUSES (mapper); + tree dummy_var = OMP_DECLARE_MAPPER_DECL (mapper); + tree mapper_name = NULL_TREE; + tree iterator = *outlist ? OMP_CLAUSE_ITERATORS (*outlist) : NULL_TREE; + + remap_mapper_decl_info map_info; + map_info.dummy_var = dummy_var; + map_info.expr = expr; + + for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) + { + tree unshared = unshare_expr (c); + enum gomp_map_kind clause_kind = OMP_CLAUSE_MAP_KIND (c); + tree t = OMP_CLAUSE_DECL (unshared); + tree type = NULL_TREE; + bool nonunit_array_with_mapper = false; + + if (clause_kind == GOMP_MAP_PUSH_MAPPER_NAME) + { + mapper_name = t; + continue; + } + else if (clause_kind == GOMP_MAP_POP_MAPPER_NAME) + { + mapper_name = NULL_TREE; + continue; + } + + if (TREE_CODE (t) == OMP_ARRAY_SECTION) + { + tree t2 = lang_hooks.decls.omp_map_array_section (loc, t); + + if (t2 == t) + { + nonunit_array_with_mapper = true; + /* We'd want use the mapper for the element type if this worked: + look that one up. */ + type = TREE_TYPE (TREE_TYPE (t)); + } + else + { + t = t2; + type = TREE_TYPE (t); + } + } + else + type = TREE_TYPE (t); + + gcc_assert (type); + + if (type == error_mark_node) + continue; + + walk_tree (&unshared, remap_mapper_decl_1, &map_info, NULL); + + OMP_CLAUSE_LOCATION (unshared) = loc; + + enum gomp_map_kind decayed_kind + = omp_map_decayed_kind (clause_kind, outer_kind, + (ort & C_ORT_EXIT_DATA) != 0 + || (outer_kind == GOMP_MAP_FROM + && (ort & C_ORT_UPDATE) != 0)); + OMP_CLAUSE_SET_MAP_KIND (unshared, decayed_kind); + + type = TYPE_MAIN_VARIANT (type); + + tree mapper_fn = lang_hooks.decls.omp_mapper_lookup (mapper_name, type); + + if (mapper_fn && nonunit_array_with_mapper) + { + sorry ("user-defined mapper with non-unit length array section"); + continue; + } + else if (mapper_fn) + { + tree nested_mapper + = lang_hooks.decls.omp_extract_mapper_directive (mapper_fn); + if (nested_mapper != mapper) + { + outlist = omp_instantiate_mapper (loc, outlist, nested_mapper, + t, outer_kind, ort); + continue; + } + } + else if (mapper_name) + { + error ("mapper %qE not found for type %qT", mapper_name, type); + continue; + } + + if (ort & C_ORT_UPDATE) + { + bool force_p, always_p, present_p; + decayed_kind + = omp_split_map_kind (decayed_kind, &force_p, &always_p, + &present_p); + /* We don't expect to see these flags here. */ + gcc_assert (!force_p && !always_p); + /* For a "target update" operation, we want to turn the map node + expanded from the mapper back into a OMP_CLAUSE_TO or + OMP_CLAUSE_FROM node. If we can do neither, emit a warning and + drop the clause. */ + switch (decayed_kind) + { + case GOMP_MAP_TO: + case GOMP_MAP_FROM: + { + tree xfer + = build_omp_clause (loc, (decayed_kind == GOMP_MAP_TO + ? OMP_CLAUSE_TO : OMP_CLAUSE_FROM)); + OMP_CLAUSE_DECL (xfer) = OMP_CLAUSE_DECL (unshared); + OMP_CLAUSE_SIZE (xfer) = OMP_CLAUSE_SIZE (unshared); + /* For FROM/TO clauses, "present" is represented by a flag. + Set it for the expanded clause here. */ + if (present_p) + OMP_CLAUSE_MOTION_PRESENT (xfer) = 1; + *outlist = xfer; + outlist = &OMP_CLAUSE_CHAIN (xfer); + } + break; + default: + clause_kind + = omp_split_map_kind (clause_kind, &force_p, &always_p, + &present_p); + warning_at (loc, 0, "dropping %qs clause during mapper expansion " + "in %<#pragma omp target update%>", + omp_basic_map_kind_name (clause_kind)); + inform (OMP_CLAUSE_LOCATION (c), "for map clause here"); + } + } + else + { + OMP_CLAUSE_ITERATORS (unshared) = iterator; + *outlist = unshared; + outlist = &OMP_CLAUSE_CHAIN (unshared); + } + } + + return outlist; +} + +/* Given a list of CLAUSES, scan each clause and invoke a user-defined mapper + appropriate to the type of the data in that clause, if such a mapper is + visible in the current parsing context. */ + +tree +c_omp_instantiate_mappers (tree clauses, enum c_omp_region_type ort) +{ + tree c, *pc, mapper_name = NULL_TREE; + + for (pc = &clauses, c = clauses; c; c = *pc) + { + bool using_mapper = false; + bool update_p = false, update_present_p = false; + + switch (OMP_CLAUSE_CODE (c)) + { + case OMP_CLAUSE_TO: + case OMP_CLAUSE_FROM: + update_p = true; + if (OMP_CLAUSE_MOTION_PRESENT (c)) + update_present_p = true; + /* Fallthrough. */ + case OMP_CLAUSE_MAP: + { + tree t = OMP_CLAUSE_DECL (c); + tree type = NULL_TREE; + bool nonunit_array_with_mapper = false; + + if (!update_p + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POP_MAPPER_NAME)) + { + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_PUSH_MAPPER_NAME) + mapper_name = OMP_CLAUSE_DECL (c); + else + mapper_name = NULL_TREE; + pc = &OMP_CLAUSE_CHAIN (c); + continue; + } + + if (TREE_CODE (t) == OMP_ARRAY_SECTION) + { + location_t loc = OMP_CLAUSE_LOCATION (c); + tree t2 = lang_hooks.decls.omp_map_array_section (loc, t); + + if (t2 == t) + { + /* !!! Array sections of size >1 with mappers for elements + are hard to support. Do something here. */ + nonunit_array_with_mapper = true; + type = TREE_TYPE (TREE_TYPE (t)); + } + else + { + t = t2; + type = TREE_TYPE (t); + } + } + else + type = TREE_TYPE (t); + + if (type == NULL_TREE || type == error_mark_node) + { + pc = &OMP_CLAUSE_CHAIN (c); + continue; + } + + enum gomp_map_kind kind; + if (update_p) + { + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO) + kind = update_present_p ? GOMP_MAP_PRESENT_TO + : GOMP_MAP_TO; + else + kind = update_present_p ? GOMP_MAP_PRESENT_FROM + : GOMP_MAP_FROM; + } + else + { + kind = OMP_CLAUSE_MAP_KIND (c); + if (kind == GOMP_MAP_UNSET) + kind = GOMP_MAP_TOFROM; + } + + type = TYPE_MAIN_VARIANT (type); + + tree mapper_fn + = lang_hooks.decls.omp_mapper_lookup (mapper_name, type); + + if (mapper_fn && nonunit_array_with_mapper) + { + sorry ("user-defined mapper with non-unit length " + "array section"); + using_mapper = true; + } + else if (mapper_fn) + { + tree mapper + = lang_hooks.decls.omp_extract_mapper_directive (mapper_fn); + pc = omp_instantiate_mapper (OMP_CLAUSE_LOCATION (c), + pc, mapper, t, kind, ort); + using_mapper = true; + } + else if (mapper_name) + { + error ("mapper %qE not found for type %qT", mapper_name, type); + using_mapper = true; + } + } + break; + + default: + ; + } + + if (using_mapper) + *pc = OMP_CLAUSE_CHAIN (c); + else + pc = &OMP_CLAUSE_CHAIN (c); + } + + return clauses; } const struct c_omp_directive c_omp_directives[] = { |