diff options
31 files changed, 1766 insertions, 143 deletions
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index af2bb5b..1ea2c72 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -16194,7 +16194,8 @@ c_parser_omp_clause_map (c_parser *parser, tree list) c_parser_consume_token (parser); } - nl = c_parser_omp_variable_list (parser, clause_loc, OMP_CLAUSE_MAP, list); + nl = c_parser_omp_variable_list (parser, clause_loc, OMP_CLAUSE_MAP, list, + true); for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c)) OMP_CLAUSE_SET_MAP_KIND (c, kind); diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index 7524304..b77ecb6 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -13241,6 +13241,11 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types, return error_mark_node; } t = TREE_OPERAND (t, 0); + if (TREE_CODE (t) == MEM_REF) + { + t = TREE_OPERAND (t, 0); + STRIP_NOPS (t); + } if (ort == C_ORT_ACC && TREE_CODE (t) == MEM_REF) { if (maybe_ne (mem_ref_offset (t), 0)) @@ -14085,6 +14090,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) tree ordered_clause = NULL_TREE; tree schedule_clause = NULL_TREE; bool oacc_async = false; + bool indir_component_ref_p = false; tree last_iterators = NULL_TREE; bool last_iterators_remove = false; tree *nogroup_seen = NULL; @@ -14886,6 +14892,11 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) { while (TREE_CODE (t) == COMPONENT_REF) t = TREE_OPERAND (t, 0); + if (TREE_CODE (t) == MEM_REF) + { + t = TREE_OPERAND (t, 0); + STRIP_NOPS (t); + } if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_IMPLICIT (c) && (bitmap_bit_p (&map_head, DECL_UID (t)) @@ -14952,6 +14963,14 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) bias) to zero here, so it is not set erroneously to the pointer size later on in gimplify.c. */ OMP_CLAUSE_SIZE (c) = size_zero_node; + indir_component_ref_p = false; + if (TREE_CODE (t) == COMPONENT_REF + && TREE_CODE (TREE_OPERAND (t, 0)) == MEM_REF) + { + t = TREE_OPERAND (TREE_OPERAND (t, 0), 0); + indir_component_ref_p = true; + STRIP_NOPS (t); + } if (TREE_CODE (t) == COMPONENT_REF && OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_) { @@ -15024,6 +15043,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) else if ((OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP || (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER)) + && !indir_component_ref_p && !c_mark_addressable (t)) remove = true; else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP @@ -15080,8 +15100,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) bitmap_set_bit (&map_firstprivate_head, DECL_UID (t)); } else if (bitmap_bit_p (&map_head, DECL_UID (t)) - && (ort == C_ORT_ACC - || !bitmap_bit_p (&map_field_head, DECL_UID (t)))) + && !bitmap_bit_p (&map_field_head, DECL_UID (t))) { if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) error_at (OMP_CLAUSE_LOCATION (c), diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h index 3510512..b1c3bc5 100644 --- a/gcc/cp/cp-tree.h +++ b/gcc/cp/cp-tree.h @@ -7671,6 +7671,8 @@ extern tree start_lambda_function (tree fn, tree lambda_expr); extern void finish_lambda_function (tree body); extern bool regenerated_lambda_fn_p (tree); extern tree most_general_lambda (tree); +extern tree finish_omp_target (location_t, tree, tree, bool); +extern void finish_omp_target_clauses (location_t, tree, tree *); /* in tree.c */ extern int cp_tree_operand_length (const_tree); diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index 55e6a1a..ed36709 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -39315,7 +39315,7 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list) } nlist = cp_parser_omp_var_list_no_open (parser, OMP_CLAUSE_MAP, list, - NULL); + NULL, true); for (c = nlist; c != list; c = OMP_CLAUSE_CHAIN (c)) OMP_CLAUSE_SET_MAP_KIND (c, kind); @@ -44105,8 +44105,6 @@ static bool cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, enum pragma_context context, bool *if_p) { - tree *pc = NULL, stmt; - if (flag_openmp) omp_requires_mask = (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_USED); @@ -44211,16 +44209,10 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, = cclauses[C_OMP_CLAUSE_SPLIT_TARGET]; cclauses[C_OMP_CLAUSE_SPLIT_TARGET] = tc; } - tree stmt = make_node (OMP_TARGET); - TREE_TYPE (stmt) = void_type_node; - OMP_TARGET_CLAUSES (stmt) = cclauses[C_OMP_CLAUSE_SPLIT_TARGET]; - c_omp_adjust_map_clauses (OMP_TARGET_CLAUSES (stmt), true); - OMP_TARGET_BODY (stmt) = body; - OMP_TARGET_COMBINED (stmt) = 1; - SET_EXPR_LOCATION (stmt, pragma_tok->location); - add_stmt (stmt); - pc = &OMP_TARGET_CLAUSES (stmt); - goto check_clauses; + c_omp_adjust_map_clauses (cclauses[C_OMP_CLAUSE_SPLIT_TARGET], true); + finish_omp_target (pragma_tok->location, + cclauses[C_OMP_CLAUSE_SPLIT_TARGET], body, true); + return true; } else if (!flag_openmp) /* flag_openmp_simd */ { @@ -44255,13 +44247,10 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, return false; } - stmt = make_node (OMP_TARGET); - TREE_TYPE (stmt) = void_type_node; - - OMP_TARGET_CLAUSES (stmt) - = cp_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK, - "#pragma omp target", pragma_tok, false); - for (tree c = OMP_TARGET_CLAUSES (stmt); c; c = OMP_CLAUSE_CHAIN (c)) + tree clauses = cp_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK, + "#pragma omp target", pragma_tok, + false); + for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IN_REDUCTION) { tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); @@ -44270,45 +44259,13 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c); OMP_CLAUSE_CHAIN (c) = nc; } - OMP_TARGET_CLAUSES (stmt) - = finish_omp_clauses (OMP_TARGET_CLAUSES (stmt), C_ORT_OMP_TARGET); - c_omp_adjust_map_clauses (OMP_TARGET_CLAUSES (stmt), true); + clauses = finish_omp_clauses (clauses, C_ORT_OMP_TARGET); - pc = &OMP_TARGET_CLAUSES (stmt); + c_omp_adjust_map_clauses (clauses, true); keep_next_level (true); - OMP_TARGET_BODY (stmt) = cp_parser_omp_structured_block (parser, if_p); - - SET_EXPR_LOCATION (stmt, pragma_tok->location); - add_stmt (stmt); + tree body = cp_parser_omp_structured_block (parser, if_p); -check_clauses: - while (*pc) - { - if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP) - switch (OMP_CLAUSE_MAP_KIND (*pc)) - { - case GOMP_MAP_TO: - case GOMP_MAP_ALWAYS_TO: - case GOMP_MAP_FROM: - case GOMP_MAP_ALWAYS_FROM: - case GOMP_MAP_TOFROM: - case GOMP_MAP_ALWAYS_TOFROM: - case GOMP_MAP_ALLOC: - case GOMP_MAP_FIRSTPRIVATE_POINTER: - case GOMP_MAP_FIRSTPRIVATE_REFERENCE: - case GOMP_MAP_ALWAYS_POINTER: - case GOMP_MAP_ATTACH_DETACH: - break; - default: - error_at (OMP_CLAUSE_LOCATION (*pc), - "%<#pragma omp target%> with map-type other " - "than %<to%>, %<from%>, %<tofrom%> or %<alloc%> " - "on %<map%> clause"); - *pc = OMP_CLAUSE_CHAIN (*pc); - continue; - } - pc = &OMP_CLAUSE_CHAIN (*pc); - } + finish_omp_target (pragma_tok->location, clauses, body, false); return true; } diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c index d3efc6e..9834baf 100644 --- a/gcc/cp/pt.c +++ b/gcc/cp/pt.c @@ -18975,6 +18975,11 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl, t = copy_node (t); OMP_BODY (t) = stmt; OMP_CLAUSES (t) = tmp; + + if (TREE_CODE (t) == OMP_TARGET) + finish_omp_target_clauses (EXPR_LOCATION (t), OMP_BODY (t), + &OMP_CLAUSES (t)); + if (TREE_CODE (t) == OMP_TARGET && OMP_TARGET_COMBINED (t)) { tree teams = cp_walk_tree (&stmt, tsubst_find_omp_teams, NULL, NULL); diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index edba4b6..9ebb1ac 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -5049,15 +5049,16 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types, return error_mark_node; } t = TREE_OPERAND (t, 0); - if (ort == C_ORT_ACC && TREE_CODE (t) == INDIRECT_REF) - t = TREE_OPERAND (t, 0); + if (TREE_CODE (t) == INDIRECT_REF) + { + t = TREE_OPERAND (t, 0); + STRIP_NOPS (t); + } } if (REFERENCE_REF_P (t)) t = TREE_OPERAND (t, 0); } - if (TREE_CODE (t) == FIELD_DECL - && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_AFFINITY - || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND)) + if (TREE_CODE (t) == FIELD_DECL) ret = finish_non_static_data_member (t, NULL_TREE, NULL_TREE); else if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL) { @@ -5073,18 +5074,6 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types, omp_clause_code_name[OMP_CLAUSE_CODE (c)]); return error_mark_node; } - else if ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP - && TREE_CODE (t) == PARM_DECL - && DECL_ARTIFICIAL (t) - && DECL_NAME (t) == this_identifier - && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_AFFINITY - && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND) - { - error_at (OMP_CLAUSE_LOCATION (c), - "%<this%> allowed in OpenMP only in %<declare simd%>" - " clauses"); - return error_mark_node; - } else if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_AFFINITY && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND && VAR_P (t) && CP_DECL_THREAD_LOCAL_P (t)) @@ -5599,6 +5588,8 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort) } OMP_CLAUSE_DECL (c) = first; OMP_CLAUSE_SIZE (c) = size; + if (TREE_CODE (t) == FIELD_DECL) + t = finish_non_static_data_member (t, NULL_TREE, NULL_TREE); if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP || (TREE_CODE (t) == COMPONENT_REF && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE)) @@ -6611,6 +6602,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) bool order_seen = false; bool schedule_seen = false; bool oacc_async = false; + bool indir_component_ref_p = false; tree last_iterators = NULL_TREE; bool last_iterators_remove = false; /* 1 if normal/task reduction has been seen, -1 if inscan reduction @@ -7862,6 +7854,11 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) t = TREE_OPERAND (t, 0); if (REFERENCE_REF_P (t)) t = TREE_OPERAND (t, 0); + if (TREE_CODE (t) == INDIRECT_REF) + { + t = TREE_OPERAND (t, 0); + STRIP_NOPS (t); + } if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_IMPLICIT (c) && (bitmap_bit_p (&map_head, DECL_UID (t)) @@ -7934,9 +7931,14 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) t = TREE_OPERAND (t, 0); OMP_CLAUSE_DECL (c) = t; } + indir_component_ref_p = false; if (TREE_CODE (t) == COMPONENT_REF && TREE_CODE (TREE_OPERAND (t, 0)) == INDIRECT_REF) - t = TREE_OPERAND (TREE_OPERAND (t, 0), 0); + { + t = TREE_OPERAND (TREE_OPERAND (t, 0), 0); + indir_component_ref_p = true; + STRIP_NOPS (t); + } if (TREE_CODE (t) == COMPONENT_REF && OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_) { @@ -7983,6 +7985,13 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) goto handle_map_references; } } + if (!processing_template_decl + && TREE_CODE (t) == FIELD_DECL) + { + OMP_CLAUSE_DECL (c) = finish_non_static_data_member (t, NULL_TREE, + NULL_TREE); + break; + } if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL) { if (processing_template_decl && TREE_CODE (t) != OVERLOAD) @@ -8009,19 +8018,12 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) omp_clause_code_name[OMP_CLAUSE_CODE (c)]); remove = true; } - else if (ort != C_ORT_ACC && t == current_class_ptr) - { - error_at (OMP_CLAUSE_LOCATION (c), - "%<this%> allowed in OpenMP only in %<declare simd%>" - " clauses"); - remove = true; - break; - } else if (!processing_template_decl && !TYPE_REF_P (TREE_TYPE (t)) && (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP || (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER)) + && !indir_component_ref_p && !cxx_mark_addressable (t)) remove = true; else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP @@ -9177,6 +9179,533 @@ finish_omp_construct (enum tree_code code, tree body, tree clauses) return add_stmt (stmt); } +/* Used to walk OpenMP target directive body. */ + +struct omp_target_walk_data +{ + /* Holds the 'this' expression found in current function. */ + tree current_object; + + /* True if the 'this' expression was accessed in the target body. */ + bool this_expr_accessed; + + /* For non-static functions, record which pointer-typed members were + accessed, and the whole expression. */ + hash_map<tree, tree> ptr_members_accessed; + + /* Record which lambda objects were accessed in target body. */ + hash_set<tree> lambda_objects_accessed; + + /* For lambda functions, the __closure object expression of the current + function, and the set of captured variables accessed in target body. */ + tree current_closure; + hash_set<tree> closure_vars_accessed; + + /* Local variables declared inside a BIND_EXPR, used to filter out such + variables when recording lambda_objects_accessed. */ + hash_set<tree> local_decls; +}; + +/* Helper function of finish_omp_target_clauses, called via + cp_walk_tree_without_duplicates. Traverse body of OpenMP target + directive *TP, and fill out omp_target_walk_data passed in *PTR. */ + +static tree +finish_omp_target_clauses_r (tree *tp, int *walk_subtrees, void *ptr) +{ + tree t = *tp; + struct omp_target_walk_data *data = (struct omp_target_walk_data *) ptr; + tree current_object = data->current_object; + tree current_closure = data->current_closure; + + /* References inside of these expression codes shouldn't incur any + form of mapping, so return early. */ + if (TREE_CODE (t) == SIZEOF_EXPR + || TREE_CODE (t) == ALIGNOF_EXPR) + { + *walk_subtrees = 0; + return NULL_TREE; + } + + if (current_object) + { + tree this_expr = TREE_OPERAND (current_object, 0); + + if (operand_equal_p (t, this_expr)) + { + data->this_expr_accessed = true; + *walk_subtrees = 0; + return NULL_TREE; + } + + if (TREE_CODE (t) == COMPONENT_REF + && POINTER_TYPE_P (TREE_TYPE (t)) + && operand_equal_p (TREE_OPERAND (t, 0), current_object) + && TREE_CODE (TREE_OPERAND (t, 1)) == FIELD_DECL) + { + data->this_expr_accessed = true; + tree fld = TREE_OPERAND (t, 1); + if (data->ptr_members_accessed.get (fld) == NULL) + { + if (TREE_CODE (TREE_TYPE (t)) == REFERENCE_TYPE) + t = convert_from_reference (t); + data->ptr_members_accessed.put (fld, t); + } + *walk_subtrees = 0; + return NULL_TREE; + } + } + + /* When the current_function_decl is a lambda function, the closure object + argument's type seems to not yet have fields layed out, so a recording + of DECL_VALUE_EXPRs during the target body walk seems the only way to + find them. */ + if (current_closure + && (TREE_CODE (t) == VAR_DECL + || TREE_CODE (t) == PARM_DECL + || TREE_CODE (t) == RESULT_DECL) + && DECL_HAS_VALUE_EXPR_P (t) + && TREE_CODE (DECL_VALUE_EXPR (t)) == COMPONENT_REF + && operand_equal_p (current_closure, + TREE_OPERAND (DECL_VALUE_EXPR (t), 0))) + { + if (!data->closure_vars_accessed.contains (t)) + data->closure_vars_accessed.add (t); + *walk_subtrees = 0; + return NULL_TREE; + } + + if (TREE_CODE (t) == BIND_EXPR) + { + tree block = BIND_EXPR_BLOCK (t); + for (tree var = BLOCK_VARS (block); var; var = DECL_CHAIN (var)) + if (!data->local_decls.contains (var)) + data->local_decls.add (var); + return NULL_TREE; + } + + if (TREE_TYPE (t) && LAMBDA_TYPE_P (TREE_TYPE (t))) + { + tree lt = TREE_TYPE (t); + gcc_assert (CLASS_TYPE_P (lt)); + + if (!data->lambda_objects_accessed.contains (t) + /* Do not prepare to create target maps for locally declared + lambdas or anonymous ones. */ + && !data->local_decls.contains (t) + && TREE_CODE (t) != TARGET_EXPR) + data->lambda_objects_accessed.add (t); + *walk_subtrees = 0; + return NULL_TREE; + } + + return NULL_TREE; +} + +/* Helper function for finish_omp_target, and also from tsubst_expr. + Create additional clauses for mapping of non-static members, lambda objects, + etc. */ + +void +finish_omp_target_clauses (location_t loc, tree body, tree *clauses_ptr) +{ + omp_target_walk_data data; + data.this_expr_accessed = false; + + tree ct = current_nonlambda_class_type (); + if (ct) + { + tree object = maybe_dummy_object (ct, NULL); + object = maybe_resolve_dummy (object, true); + data.current_object = object; + } + else + data.current_object = NULL_TREE; + + if (DECL_LAMBDA_FUNCTION_P (current_function_decl)) + { + tree closure = DECL_ARGUMENTS (current_function_decl); + data.current_closure = build_indirect_ref (loc, closure, RO_UNARY_STAR); + } + else + data.current_closure = NULL_TREE; + + cp_walk_tree_without_duplicates (&body, finish_omp_target_clauses_r, &data); + + auto_vec<tree, 16> new_clauses; + + tree omp_target_this_expr = NULL_TREE; + tree *explicit_this_deref_map = NULL; + if (data.this_expr_accessed) + { + omp_target_this_expr = TREE_OPERAND (data.current_object, 0); + + /* See if explicit user-specified map(this[:]) clause already exists. + If not, we create an implicit map(tofrom:this[:1]) clause. */ + for (tree *cp = clauses_ptr; *cp; cp = &OMP_CLAUSE_CHAIN (*cp)) + if (OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP + && (TREE_CODE (OMP_CLAUSE_DECL (*cp)) == INDIRECT_REF + || TREE_CODE (OMP_CLAUSE_DECL (*cp)) == MEM_REF) + && operand_equal_p (TREE_OPERAND (OMP_CLAUSE_DECL (*cp), 0), + omp_target_this_expr)) + { + explicit_this_deref_map = cp; + break; + } + } + + if (DECL_LAMBDA_FUNCTION_P (current_function_decl) + && (data.this_expr_accessed + || !data.closure_vars_accessed.is_empty ())) + { + /* For lambda functions, we need to first create a copy of the + __closure object. */ + tree closure = DECL_ARGUMENTS (current_function_decl); + tree c = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TO); + OMP_CLAUSE_DECL (c) + = build_indirect_ref (loc, closure, RO_UNARY_STAR); + OMP_CLAUSE_SIZE (c) + = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (closure))); + new_clauses.safe_push (c); + + tree closure_obj = OMP_CLAUSE_DECL (c); + tree closure_type = TREE_TYPE (closure_obj); + + gcc_assert (LAMBDA_TYPE_P (closure_type) + && CLASS_TYPE_P (closure_type)); + + tree c2 = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER); + OMP_CLAUSE_DECL (c2) = closure; + OMP_CLAUSE_SIZE (c2) = size_zero_node; + new_clauses.safe_push (c2); + } + + if (data.this_expr_accessed) + { + /* If the this-expr was accessed, create a map(*this) clause. */ + enum gomp_map_kind kind = GOMP_MAP_TOFROM; + if (explicit_this_deref_map) + { + tree this_map = *explicit_this_deref_map; + tree nc = OMP_CLAUSE_CHAIN (this_map); + gcc_assert (nc != NULL_TREE + && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (nc) + == GOMP_MAP_FIRSTPRIVATE_POINTER)); + kind = OMP_CLAUSE_MAP_KIND (this_map); + /* Remove the original 'map(*this) map(firstprivate_ptr:this)' + two-map sequence away from the chain. */ + *explicit_this_deref_map = OMP_CLAUSE_CHAIN (nc); + } + tree c = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c, kind); + OMP_CLAUSE_DECL (c) + = build_indirect_ref (loc, omp_target_this_expr, RO_UNARY_STAR); + OMP_CLAUSE_SIZE (c) + = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (omp_target_this_expr))); + new_clauses.safe_push (c); + + /* If we're in a lambda function, the this-pointer will actually be + '__closure->this', a mapped member of __closure, hence always_pointer. + Otherwise it's a firstprivate pointer. */ + enum gomp_map_kind ptr_kind + = (DECL_LAMBDA_FUNCTION_P (current_function_decl) + ? GOMP_MAP_ALWAYS_POINTER + : GOMP_MAP_FIRSTPRIVATE_POINTER); + c = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c, ptr_kind); + OMP_CLAUSE_DECL (c) = omp_target_this_expr; + OMP_CLAUSE_SIZE (c) = size_zero_node; + new_clauses.safe_push (c); + } + + if (DECL_LAMBDA_FUNCTION_P (current_function_decl)) + { + if (omp_target_this_expr) + { + STRIP_NOPS (omp_target_this_expr); + gcc_assert (DECL_HAS_VALUE_EXPR_P (omp_target_this_expr)); + omp_target_this_expr = DECL_VALUE_EXPR (omp_target_this_expr); + } + + for (hash_set<tree>::iterator i = data.closure_vars_accessed.begin (); + i != data.closure_vars_accessed.end (); ++i) + { + tree orig_decl = *i; + tree closure_expr = DECL_VALUE_EXPR (orig_decl); + + if (TREE_CODE (TREE_TYPE (orig_decl)) == POINTER_TYPE + || TREE_CODE (TREE_TYPE (orig_decl)) == REFERENCE_TYPE) + { + /* this-pointer is processed above, outside this loop. */ + if (omp_target_this_expr + && operand_equal_p (closure_expr, omp_target_this_expr)) + continue; + + bool ptr_p = TREE_CODE (TREE_TYPE (orig_decl)) == POINTER_TYPE; + enum gomp_map_kind kind, ptr_kind, nc_kind; + tree size; + + if (ptr_p) + { + /* For pointers, default mapped as zero-length array + section. */ + kind = GOMP_MAP_ALLOC; + nc_kind = GOMP_MAP_FIRSTPRIVATE_POINTER; + ptr_kind = GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION; + size = size_zero_node; + } + else + { + /* For references, default mapped as appearing on map + clause. */ + kind = GOMP_MAP_TOFROM; + nc_kind = GOMP_MAP_FIRSTPRIVATE_REFERENCE; + ptr_kind = GOMP_MAP_ALWAYS_POINTER; + size = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (closure_expr))); + } + + for (tree *p = clauses_ptr; *p; p = &OMP_CLAUSE_CHAIN (*p)) + if (OMP_CLAUSE_CODE (*p) == OMP_CLAUSE_MAP + && (TREE_CODE (OMP_CLAUSE_DECL (*p)) == INDIRECT_REF + || TREE_CODE (OMP_CLAUSE_DECL (*p)) == MEM_REF) + && operand_equal_p (TREE_OPERAND (OMP_CLAUSE_DECL (*p), 0), + orig_decl)) + { + /* If this was already specified by user as a map, + save the user specified map kind, delete the + "map(*ptr/ref), map(firstprivate ptr/ref)" sequence, + and insert our own sequence: + "map(*__closure->ptr/ref), map(<ptr_kind>:__closure->ref" + */ + tree nc = OMP_CLAUSE_CHAIN (*p); + gcc_assert (nc != NULL_TREE + && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (nc) == nc_kind); + /* Update with user specified kind and size. */ + kind = OMP_CLAUSE_MAP_KIND (*p); + size = OMP_CLAUSE_SIZE (*p); + *p = OMP_CLAUSE_CHAIN (nc); + break; + } + + tree c = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c, kind); + OMP_CLAUSE_DECL (c) + = build_indirect_ref (loc, closure_expr, RO_UNARY_STAR); + OMP_CLAUSE_SIZE (c) = size; + if (ptr_p) + OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1; + new_clauses.safe_push (c); + + c = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c, ptr_kind); + OMP_CLAUSE_DECL (c) = closure_expr; + OMP_CLAUSE_SIZE (c) = size_zero_node; + new_clauses.safe_push (c); + } + } + } + + if (!data.ptr_members_accessed.is_empty ()) + for (hash_map<tree, tree>::iterator i = data.ptr_members_accessed.begin (); + i != data.ptr_members_accessed.end (); ++i) + { + /* For each referenced member that is of pointer or reference-to-pointer + type, create the equivalent of map(alloc:this->ptr[:0]). */ + tree field_decl = (*i).first; + tree ptr_member = (*i).second; + + for (tree c = *clauses_ptr; c; c = OMP_CLAUSE_CHAIN (c)) + { + /* If map(this->ptr[:N] already exists, avoid creating another + such map. */ + tree decl = OMP_CLAUSE_DECL (c); + if ((TREE_CODE (decl) == INDIRECT_REF + || TREE_CODE (decl) == MEM_REF) + && operand_equal_p (TREE_OPERAND (decl, 0), ptr_member)) + goto next_ptr_member; + } + + if (!cxx_mark_addressable (ptr_member)) + gcc_unreachable (); + + if (TREE_CODE (TREE_TYPE (field_decl)) == REFERENCE_TYPE) + { + /* For reference to pointers, we need to map the referenced + pointer first for things to be correct. */ + tree ptr_member_type = TREE_TYPE (ptr_member); + + /* Map pointer target as zero-length array section. */ + tree c = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALLOC); + OMP_CLAUSE_DECL (c) + = build1 (INDIRECT_REF, TREE_TYPE (ptr_member_type), ptr_member); + OMP_CLAUSE_SIZE (c) = size_zero_node; + OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1; + + /* Map pointer to zero-length array section. */ + tree c2 = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND + (c2, GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION); + OMP_CLAUSE_DECL (c2) = ptr_member; + OMP_CLAUSE_SIZE (c2) = size_zero_node; + + /* Attach reference-to-pointer field to pointer. */ + tree c3 = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c3, GOMP_MAP_ATTACH); + OMP_CLAUSE_DECL (c3) = TREE_OPERAND (ptr_member, 0); + OMP_CLAUSE_SIZE (c3) = size_zero_node; + + new_clauses.safe_push (c); + new_clauses.safe_push (c2); + new_clauses.safe_push (c3); + } + else if (TREE_CODE (TREE_TYPE (field_decl)) == POINTER_TYPE) + { + /* Map pointer target as zero-length array section. */ + tree c = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALLOC); + OMP_CLAUSE_DECL (c) = build_indirect_ref (loc, ptr_member, + RO_UNARY_STAR); + OMP_CLAUSE_SIZE (c) = size_zero_node; + OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1; + + /* Attach zero-length array section to pointer. */ + tree c2 = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND + (c2, GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION); + OMP_CLAUSE_DECL (c2) = ptr_member; + OMP_CLAUSE_SIZE (c2) = size_zero_node; + + new_clauses.safe_push (c); + new_clauses.safe_push (c2); + } + else + gcc_unreachable (); + + next_ptr_member: + ; + } + + for (hash_set<tree>::iterator i = data.lambda_objects_accessed.begin (); + i != data.lambda_objects_accessed.end (); ++i) + { + tree lobj = *i; + if (TREE_CODE (lobj) == TARGET_EXPR) + lobj = TREE_OPERAND (lobj, 0); + + tree lt = TREE_TYPE (lobj); + gcc_assert (LAMBDA_TYPE_P (lt) && CLASS_TYPE_P (lt)); + + tree lc = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (lc, GOMP_MAP_TO); + OMP_CLAUSE_DECL (lc) = lobj; + OMP_CLAUSE_SIZE (lc) = TYPE_SIZE_UNIT (lt); + new_clauses.safe_push (lc); + + for (tree fld = TYPE_FIELDS (lt); fld; fld = DECL_CHAIN (fld)) + { + if (TREE_CODE (TREE_TYPE (fld)) == POINTER_TYPE) + { + tree exp = build3 (COMPONENT_REF, TREE_TYPE (fld), + lobj, fld, NULL_TREE); + tree c = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALLOC); + OMP_CLAUSE_DECL (c) + = build_indirect_ref (loc, exp, RO_UNARY_STAR); + OMP_CLAUSE_SIZE (c) = size_zero_node; + OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1; + new_clauses.safe_push (c); + + c = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND + (c, GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION); + OMP_CLAUSE_DECL (c) = exp; + OMP_CLAUSE_SIZE (c) = size_zero_node; + new_clauses.safe_push (c); + } + else if (TREE_CODE (TREE_TYPE (fld)) == REFERENCE_TYPE) + { + tree exp = build3 (COMPONENT_REF, TREE_TYPE (fld), + lobj, fld, NULL_TREE); + tree c = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TOFROM); + OMP_CLAUSE_DECL (c) + = build1 (INDIRECT_REF, TREE_TYPE (TREE_TYPE (exp)), exp); + OMP_CLAUSE_SIZE (c) + = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (exp))); + new_clauses.safe_push (c); + + c = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER); + OMP_CLAUSE_DECL (c) = exp; + OMP_CLAUSE_SIZE (c) = size_zero_node; + new_clauses.safe_push (c); + } + } + } + + tree c = *clauses_ptr; + for (int i = new_clauses.length () - 1; i >= 0; i--) + { + OMP_CLAUSE_CHAIN (new_clauses[i]) = c; + c = new_clauses[i]; + } + *clauses_ptr = c; +} + +/* Called from cp_parser_omp_target. Create additional implicit clauses for + OpenMP target directives, and do sanity checks. */ + +tree +finish_omp_target (location_t loc, tree clauses, tree body, bool combined_p) +{ + if (!processing_template_decl) + finish_omp_target_clauses (loc, body, &clauses); + + tree stmt = make_node (OMP_TARGET); + TREE_TYPE (stmt) = void_type_node; + OMP_TARGET_CLAUSES (stmt) = clauses; + OMP_TARGET_BODY (stmt) = body; + OMP_TARGET_COMBINED (stmt) = combined_p; + SET_EXPR_LOCATION (stmt, loc); + + tree c = clauses; + while (c) + { + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP) + switch (OMP_CLAUSE_MAP_KIND (c)) + { + case GOMP_MAP_TO: + case GOMP_MAP_ALWAYS_TO: + case GOMP_MAP_FROM: + case GOMP_MAP_ALWAYS_FROM: + case GOMP_MAP_TOFROM: + case GOMP_MAP_ALWAYS_TOFROM: + case GOMP_MAP_ALLOC: + case GOMP_MAP_FIRSTPRIVATE_POINTER: + case GOMP_MAP_FIRSTPRIVATE_REFERENCE: + case GOMP_MAP_ALWAYS_POINTER: + case GOMP_MAP_ATTACH_DETACH: + case GOMP_MAP_ATTACH: + case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION: + case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION: + break; + default: + error_at (OMP_CLAUSE_LOCATION (c), + "%<#pragma omp target%> with map-type other " + "than %<to%>, %<from%>, %<tofrom%> or %<alloc%> " + "on %<map%> clause"); + break; + } + c = OMP_CLAUSE_CHAIN (c); + } + return add_stmt (stmt); +} + tree finish_omp_parallel (tree clauses, tree body) { diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 326476f..0a55bf7 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -53,6 +53,7 @@ along with GCC; see the file COPYING3. If not see #include "langhooks.h" #include "tree-cfg.h" #include "tree-ssa.h" +#include "tree-hash-traits.h" #include "omp-general.h" #include "omp-low.h" #include "gimple-low.h" @@ -8921,7 +8922,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, { struct gimplify_omp_ctx *ctx, *outer_ctx; tree c; - hash_map<tree, tree> *struct_map_to_clause = NULL; + hash_map<tree_operand_hash, tree> *struct_map_to_clause = NULL; hash_set<tree> *struct_deref_set = NULL; tree *prev_list_p = NULL, *orig_list_p = list_p; int handled_depend_iterators = -1; @@ -9365,7 +9366,14 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, GOVD_FIRSTPRIVATE | GOVD_SEEN); } - if (!DECL_P (decl)) + if (TREE_CODE (decl) == TARGET_EXPR) + { + if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p, NULL, + is_gimple_lvalue, fb_lvalue) + == GS_ERROR) + remove = true; + } + else if (!DECL_P (decl)) { tree d = decl, *pd; if (TREE_CODE (d) == ARRAY_REF) @@ -9381,12 +9389,15 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, && TREE_CODE (decl) == INDIRECT_REF && TREE_CODE (TREE_OPERAND (decl, 0)) == COMPONENT_REF && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0))) - == REFERENCE_TYPE)) + == REFERENCE_TYPE) + && (OMP_CLAUSE_MAP_KIND (c) + != GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION)) { pd = &TREE_OPERAND (decl, 0); decl = TREE_OPERAND (decl, 0); } bool indir_p = false; + bool component_ref_p = false; tree orig_decl = decl; tree decl_ref = NULL_TREE; if ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA)) != 0 @@ -9397,6 +9408,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, while (TREE_CODE (decl) == COMPONENT_REF) { decl = TREE_OPERAND (decl, 0); + component_ref_p = true; if (((TREE_CODE (decl) == MEM_REF && integer_zerop (TREE_OPERAND (decl, 1))) || INDIRECT_REF_P (decl)) @@ -9405,6 +9417,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, { indir_p = true; decl = TREE_OPERAND (decl, 0); + STRIP_NOPS (decl); } if (TREE_CODE (decl) == INDIRECT_REF && DECL_P (TREE_OPERAND (decl, 0)) @@ -9416,8 +9429,11 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, } } } - else if (TREE_CODE (decl) == COMPONENT_REF) + else if (TREE_CODE (decl) == COMPONENT_REF + && (OMP_CLAUSE_MAP_KIND (c) + != GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION)) { + component_ref_p = true; while (TREE_CODE (decl) == COMPONENT_REF) decl = TREE_OPERAND (decl, 0); if (TREE_CODE (decl) == INDIRECT_REF @@ -9487,7 +9503,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, if (code == OACC_UPDATE && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER); - if (DECL_P (decl) + if ((DECL_P (decl) + || (component_ref_p + && (INDIRECT_REF_P (decl) + || TREE_CODE (decl) == MEM_REF))) && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH @@ -9544,7 +9563,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, gcc_assert (base == decl); splay_tree_node n - = splay_tree_lookup (ctx->variables, (splay_tree_key)decl); + = (DECL_P (decl) + ? splay_tree_lookup (ctx->variables, + (splay_tree_key) decl) + : NULL); bool ptr = (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER); bool attach_detach = (OMP_CLAUSE_MAP_KIND (c) @@ -9570,7 +9592,11 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, OMP_CLAUSE_SET_MAP_KIND (c, k); has_attachments = true; } - if (n == NULL || (n->value & GOVD_MAP) == 0) + if ((DECL_P (decl) + && (n == NULL || (n->value & GOVD_MAP) == 0)) + || (!DECL_P (decl) + && (!struct_map_to_clause + || struct_map_to_clause->get (decl) == NULL))) { tree l = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); @@ -9581,7 +9607,18 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, if (base_ref) OMP_CLAUSE_DECL (l) = unshare_expr (base_ref); else - OMP_CLAUSE_DECL (l) = decl; + { + OMP_CLAUSE_DECL (l) = unshare_expr (decl); + if (!DECL_P (OMP_CLAUSE_DECL (l)) + && (gimplify_expr (&OMP_CLAUSE_DECL (l), + pre_p, NULL, is_gimple_lvalue, + fb_lvalue) + == GS_ERROR)) + { + remove = true; + break; + } + } OMP_CLAUSE_SIZE (l) = (!attach ? size_int (1) @@ -9589,7 +9626,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, ? DECL_SIZE_UNIT (OMP_CLAUSE_DECL (l)) : TYPE_SIZE_UNIT (TREE_TYPE (OMP_CLAUSE_DECL (l)))); if (struct_map_to_clause == NULL) - struct_map_to_clause = new hash_map<tree, tree>; + struct_map_to_clause + = new hash_map<tree_operand_hash, tree>; struct_map_to_clause->put (decl, l); if (ptr || attach_detach) { @@ -9623,15 +9661,41 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, flags |= GOVD_SEEN; if (has_attachments) flags |= GOVD_MAP_HAS_ATTACHMENTS; - goto do_add_decl; + + /* If this is a *pointer-to-struct expression, make sure a + firstprivate map of the base-pointer exists. */ + if (component_ref_p + && ((TREE_CODE (decl) == MEM_REF + && integer_zerop (TREE_OPERAND (decl, 1))) + || INDIRECT_REF_P (decl)) + && DECL_P (TREE_OPERAND (decl, 0)) + && !splay_tree_lookup (ctx->variables, + ((splay_tree_key) + TREE_OPERAND (decl, 0)))) + { + decl = TREE_OPERAND (decl, 0); + tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), + OMP_CLAUSE_MAP); + enum gomp_map_kind mkind + = GOMP_MAP_FIRSTPRIVATE_POINTER; + OMP_CLAUSE_SET_MAP_KIND (c2, mkind); + OMP_CLAUSE_DECL (c2) = decl; + OMP_CLAUSE_SIZE (c2) = size_zero_node; + OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c); + OMP_CLAUSE_CHAIN (c) = c2; + } + + if (DECL_P (decl)) + goto do_add_decl; } else if (struct_map_to_clause) { tree *osc = struct_map_to_clause->get (decl); tree *sc = NULL, *scp = NULL; - if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) - || ptr - || attach_detach) + if (n != NULL + && (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) + || ptr + || attach_detach)) n->value |= GOVD_SEEN; sc = &OMP_CLAUSE_CHAIN (*osc); if (*sc != c @@ -9732,6 +9796,13 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, } else if (*sc != c) { + if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, + fb_lvalue) + == GS_ERROR) + { + remove = true; + break; + } *list_p = OMP_CLAUSE_CHAIN (c); OMP_CLAUSE_CHAIN (c) = *sc; *sc = c; @@ -9867,6 +9938,24 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, break; } + /* If this was of the form map(*pointer_to_struct), then the + 'pointer_to_struct' DECL should be considered deref'ed. */ + if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALLOC + || GOMP_MAP_COPY_TO_P (OMP_CLAUSE_MAP_KIND (c)) + || GOMP_MAP_COPY_FROM_P (OMP_CLAUSE_MAP_KIND (c))) + && INDIRECT_REF_P (orig_decl) + && DECL_P (TREE_OPERAND (orig_decl, 0)) + && TREE_CODE (TREE_TYPE (orig_decl)) == RECORD_TYPE) + { + tree ptr = TREE_OPERAND (orig_decl, 0); + if (!struct_deref_set || !struct_deref_set->contains (ptr)) + { + if (!struct_deref_set) + struct_deref_set = new hash_set<tree> (); + struct_deref_set->add (ptr); + } + } + if (!remove && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH @@ -11216,6 +11305,12 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, } } } + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT + && (code == OMP_TARGET_EXIT_DATA || code == OACC_EXIT_DATA)) + { + remove = true; + break; + } if (!DECL_P (decl)) { if ((ctx->region_type & ORT_TARGET) != 0 @@ -11262,10 +11357,6 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, = OMP_CLAUSE_CHAIN (OMP_CLAUSE_CHAIN (c)); } } - else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT - && (code == OMP_TARGET_EXIT_DATA - || code == OACC_EXIT_DATA)) - remove = true; else if (DECL_SIZE (decl) && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_POINTER diff --git a/gcc/omp-low.c b/gcc/omp-low.c index de3a26e..2a07beb 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -12635,6 +12635,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) case GOMP_MAP_ALWAYS_POINTER: case GOMP_MAP_ATTACH: case GOMP_MAP_DETACH: + case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION: + case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION: break; case GOMP_MAP_IF_PRESENT: case GOMP_MAP_FORCE_ALLOC: diff --git a/gcc/testsuite/g++.dg/gomp/target-3.C b/gcc/testsuite/g++.dg/gomp/target-3.C new file mode 100644 index 0000000..f4d40ec --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/target-3.C @@ -0,0 +1,36 @@ +// { dg-do compile } +// { dg-options "-fopenmp -fdump-tree-gimple" } + +struct S +{ + int a, b; + void bar (int); +}; + +void +S::bar (int x) +{ + #pragma omp target map (alloc: a, b) + ; + #pragma omp target enter data map (alloc: a, b) +} + +template <int N> +struct T +{ + int a, b; + void bar (int); +}; + +template <int N> +void +T<N>::bar (int x) +{ + #pragma omp target map (alloc: a, b) + ; + #pragma omp target enter data map (alloc: a, b) +} + +template struct T<0>; + +/* { dg-final { scan-tree-dump-times "map\\(struct:\\*this \\\[len: 2\\\]\\) map\\(alloc:this->a \\\[len: \[0-9\]+\\\]\\) map\\(alloc:this->b \\\[len: \[0-9\]+\\\]\\)" 4 "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C new file mode 100644 index 0000000..7f83f92 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C @@ -0,0 +1,94 @@ +// We use 'auto' without a function return type, so specify dialect here +// { dg-additional-options "-std=c++14 -fdump-tree-gimple" } +#include <cstdlib> +#include <cstring> + +template <typename L> +void +omp_target_loop (int begin, int end, L loop) +{ + #pragma omp target teams distribute parallel for + for (int i = begin; i < end; i++) + loop (i); +} + +struct S +{ + int a, len; + int *ptr; + + auto merge_data_func (int *iptr, int &b) + { + auto fn = [=](void) -> bool + { + bool mapped; + #pragma omp target map(from:mapped) + { + mapped = (ptr != NULL && iptr != NULL); + if (mapped) + { + for (int i = 0; i < len; i++) + ptr[i] += a + b + iptr[i]; + } + } + return mapped; + }; + return fn; + } +}; + +int x = 1; + +int main (void) +{ + const int N = 10; + int *data1 = new int[N]; + int *data2 = new int[N]; + memset (data1, 0xab, sizeof (int) * N); + memset (data1, 0xcd, sizeof (int) * N); + + int val = 1; + int &valref = val; + #pragma omp target enter data map(alloc: data1[:N], data2[:N]) + + omp_target_loop (0, N, [=](int i) { data1[i] = val; }); + omp_target_loop (0, N, [=](int i) { data2[i] = valref + 1; }); + + #pragma omp target update from(data1[:N], data2[:N]) + + for (int i = 0; i < N; i++) + { + if (data1[i] != 1) abort (); + if (data2[i] != 2) abort (); + } + + #pragma omp target exit data map(delete: data1[:N], data2[:N]) + + int b = 8; + S s = { 4, N, data1 }; + auto f = s.merge_data_func (data2, b); + + if (f ()) abort (); + + #pragma omp target enter data map(to: data1[:N]) + if (f ()) abort (); + + #pragma omp target enter data map(to: data2[:N]) + if (!f ()) abort (); + + #pragma omp target exit data map(from: data1[:N], data2[:N]) + + for (int i = 0; i < N; i++) + { + if (data1[i] != 0xf) abort (); + if (data2[i] != 2) abort (); + } + + return 0; +} + +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer assign, bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:__closure->__iptr \[bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */ + +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(attach_zero_length_array_section:loop\.__data1 \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */ + +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(attach_zero_length_array_section:loop\.__data2 \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/target-lambda-2.C b/gcc/testsuite/g++.dg/gomp/target-lambda-2.C new file mode 100644 index 0000000..bdf2564 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/target-lambda-2.C @@ -0,0 +1,35 @@ +// We use 'auto' without a function return type, so specify dialect here +// { dg-additional-options "-std=c++14 -fdump-tree-gimple" } +#include <cstdlib> + +#define N 10 +int main (void) +{ + int X, Y; + #pragma omp target map(from: X, Y) + { + int x = 0, y = 0; + + for (int i = 0; i < N; i++) + [&] (int v) { x += v; } (i); + + auto yinc = [&y] { y++; }; + for (int i = 0; i < N; i++) + yinc (); + + X = x; + Y = y; + } + + int Xs = 0; + for (int i = 0; i < N; i++) + Xs += i; + if (X != Xs) + abort (); + + if (Y != N) + abort (); +} + +/* Make sure lambda objects do NOT appear in target maps. */ +/* { dg-final { scan-tree-dump {(?n)#pragma omp target num_teams.* map\(from:Y \[len: [0-9]+\]\) map\(from:X \[len: [0-9]+\]\)$} "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/target-this-1.C b/gcc/testsuite/g++.dg/gomp/target-this-1.C new file mode 100644 index 0000000..de93a3e --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/target-this-1.C @@ -0,0 +1,33 @@ +// { dg-do compile } +// { dg-additional-options "-fdump-tree-gimple" } +extern "C" void abort (); + +struct S +{ + int a, b, c, d; + + int sum (void) + { + int val = 0; + val += a + b + this->c + this->d; + return val; + } + + int sum_offload (void) + { + int val = 0; + #pragma omp target map(val) + val += a + b + this->c + this->d; + return val; + } +}; + +int main (void) +{ + S s = { 1, 2, 3, 4 }; + if (s.sum () != s.sum_offload ()) + abort (); + return 0; +} + +/* { dg-final { scan-tree-dump {map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/target-this-2.C b/gcc/testsuite/g++.dg/gomp/target-this-2.C new file mode 100644 index 0000000..8a76bb8 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/target-this-2.C @@ -0,0 +1,49 @@ +// We use 'auto' without a function return type, so specify dialect here +// { dg-do compile } +// { dg-additional-options "-std=c++14 -fdump-tree-gimple" } + +extern "C" void abort (); + +struct T +{ + int x, y; + + auto sum_func (int n) + { + auto fn = [=](int m) -> int + { + int v; + v = (x + y) * n + m; + return v; + }; + return fn; + } + + auto sum_func_offload (int n) + { + auto fn = [=](int m) -> int + { + int v; + #pragma omp target map(from:v) + v = (x + y) * n + m; + return v; + }; + return fn; + } + +}; + +int main (void) +{ + T a = { 1, 2 }; + + auto s1 = a.sum_func (3); + auto s2 = a.sum_func_offload (3); + + if (s1 (1) != s2 (1)) + abort (); + + return 0; +} + +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) firstprivate\(m\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:v \[len: [0-9]+\]\)} "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/target-this-3.C b/gcc/testsuite/g++.dg/gomp/target-this-3.C new file mode 100644 index 0000000..91cfbd6 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/target-this-3.C @@ -0,0 +1,105 @@ +// { dg-do compile } +// { dg-additional-options "-fdump-tree-gimple" } +#include <cstdlib> +#include <cstring> +extern "C" void abort (); + +struct S +{ + int * ptr; + int ptr_len; + + int *&refptr; + int refptr_len; + + bool set_ptr (int n) + { + bool mapped; + #pragma omp target map(from:mapped) + { + if (ptr != NULL) + for (int i = 0; i < ptr_len; i++) + ptr[i] = n; + mapped = (ptr != NULL); + } + return mapped; + } + + bool set_refptr (int n) + { + bool mapped; + #pragma omp target map(from:mapped) + { + if (refptr != NULL) + for (int i = 0; i < refptr_len; i++) + refptr[i] = n; + mapped = (refptr != NULL); + } + return mapped; + } +}; + +int main (void) +{ + #define N 10 + int *ptr1 = new int[N]; + int *ptr2 = new int[N]; + + memset (ptr1, 0, sizeof (int) * N); + memset (ptr2, 0, sizeof (int) * N); + + S s = { ptr1, N, ptr2, N }; + + bool mapped; + int val = 123; + + mapped = s.set_ptr (val); + if (mapped) + abort (); + if (s.ptr != ptr1) + abort (); + for (int i = 0; i < N; i++) + if (ptr1[i] != 0) + abort (); + + mapped = s.set_refptr (val); + if (mapped) + abort (); + if (s.refptr != ptr2) + abort (); + for (int i = 0; i < N; i++) + if (ptr2[i] != 0) + abort (); + + #pragma omp target data map(ptr1[:N]) + mapped = s.set_ptr (val); + + if (!mapped) + abort (); + if (s.set_refptr (0)) + abort (); + if (s.ptr != ptr1 || s.refptr != ptr2) + abort (); + for (int i = 0; i < N; i++) + if (ptr1[i] != val) + abort (); + + #pragma omp target data map(ptr2[:N]) + mapped = s.set_refptr (val); + + if (!mapped) + abort (); + if (s.set_ptr (0)) + abort (); + if (s.ptr != ptr1 || s.refptr != ptr2) + abort (); + for (int i = 0; i < N; i++) + if (ptr2[i] != val) + abort (); + + return 0; +} + +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) firstprivate\(n\) map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:this->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\)} "gimple" } } */ + +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) firstprivate\(n\) map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:this->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/target-this-4.C b/gcc/testsuite/g++.dg/gomp/target-this-4.C new file mode 100644 index 0000000..e4b2a71 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/target-this-4.C @@ -0,0 +1,107 @@ +// We use 'auto' without a function return type, so specify dialect here +// { dg-additional-options "-std=c++14 -fdump-tree-gimple" } +#include <cstdlib> +#include <cstring> + +struct T +{ + int *ptr; + int ptr_len; + + int *&refptr; + int refptr_len; + + auto set_ptr_func (int n) + { + auto fn = [=](void) -> bool + { + bool mapped; + #pragma omp target map(from:mapped) + { + if (ptr) + for (int i = 0; i < ptr_len; i++) + ptr[i] = n; + mapped = (ptr != NULL); + } + return mapped; + }; + return fn; + } + + auto set_refptr_func (int n) + { + auto fn = [=](void) -> bool + { + bool mapped; + #pragma omp target map(from:mapped) + { + if (refptr) + for (int i = 0; i < refptr_len; i++) + refptr[i] = n; + mapped = (refptr != NULL); + } + return mapped; + }; + return fn; + } +}; + +int main (void) +{ + #define N 10 + int *ptr1 = new int[N]; + int *ptr2 = new int[N]; + + memset (ptr1, 0, sizeof (int) * N); + memset (ptr2, 0, sizeof (int) * N); + + T a = { ptr1, N, ptr2, N }; + + auto p1 = a.set_ptr_func (1); + auto r2 = a.set_refptr_func (2); + + if (p1 ()) + abort (); + if (r2 ()) + abort (); + + if (a.ptr != ptr1) + abort (); + if (a.refptr != ptr2) + abort (); + + for (int i = 0; i < N; i++) + if (ptr1[i] != 0) + abort (); + + for (int i = 0; i < N; i++) + if (ptr2[i] != 0) + abort (); + + #pragma omp target data map(ptr1[:N], ptr2[:N]) + { + if (!p1 ()) + abort (); + if (!r2 ()) + abort (); + } + + if (a.ptr != ptr1) + abort (); + if (a.refptr != ptr2) + abort (); + + for (int i = 0; i < N; i++) + if (ptr1[i] != 1) + abort (); + + for (int i = 0; i < N; i++) + if (ptr2[i] != 2) + abort (); + + return 0; +} + +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */ + +/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:_[0-9]+->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\)} "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/target-this-5.C b/gcc/testsuite/g++.dg/gomp/target-this-5.C new file mode 100644 index 0000000..a9ac74b --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/target-this-5.C @@ -0,0 +1,34 @@ +// { dg-do compile } +// { dg-additional-options "-fdump-tree-gimple" } +extern "C" void abort (); + +template<typename T> +struct S +{ + T a, b, c, d; + + T sum (void) + { + T val = 0; + val += a + b + this->c + this->d; + return val; + } + + T sum_offload (void) + { + T val = 0; + #pragma omp target map(val) + val += a + b + this->c + this->d; + return val; + } +}; + +int main (void) +{ + S<int> s = { 1, 2, 3, 4 }; + if (s.sum () != s.sum_offload ()) + abort (); + return 0; +} + +/* { dg-final { scan-tree-dump {map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/gomp/this-2.C b/gcc/testsuite/g++.dg/gomp/this-2.C index d03b8a0..b521a4f 100644 --- a/gcc/testsuite/g++.dg/gomp/this-2.C +++ b/gcc/testsuite/g++.dg/gomp/this-2.C @@ -9,14 +9,14 @@ struct S void S::bar (int x) { - #pragma omp target map (this, x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } + #pragma omp target map (this, x) // { dg-error "cannot take the address of .this., which is an rvalue expression" } ; - #pragma omp target map (this[0], x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } + #pragma omp target map (this[0], x) ; - #pragma omp target update to (this, x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } - #pragma omp target update to (this[0], x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } - #pragma omp target update from (this, x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } - #pragma omp target update from (this[1], x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } + #pragma omp target update to (this, x) // { dg-error "cannot take the address of .this., which is an rvalue expression" } + #pragma omp target update to (this[0], x) + #pragma omp target update from (this, x) // { dg-error "cannot take the address of .this., which is an rvalue expression" } + #pragma omp target update from (this[1], x) } template <int N> @@ -29,14 +29,14 @@ template <int N> void T<N>::bar (int x) { - #pragma omp target map (this, x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } + #pragma omp target map (this, x) // { dg-error "cannot take the address of .this., which is an rvalue expression" } ; - #pragma omp target map (this[0], x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } + #pragma omp target map (this[0], x) ; - #pragma omp target update to (this, x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } - #pragma omp target update to (this[0], x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } - #pragma omp target update from (this, x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } - #pragma omp target update from (this[1], x) // { dg-error ".this. allowed in OpenMP only in .declare simd. clauses" } + #pragma omp target update to (this, x) // { dg-error "cannot take the address of .this., which is an rvalue expression" } + #pragma omp target update to (this[0], x) + #pragma omp target update from (this, x) // { dg-error "cannot take the address of .this., which is an rvalue expression" } + #pragma omp target update from (this[1], x) } template struct T<0>; diff --git a/gcc/testsuite/gcc.dg/gomp/target-3.c b/gcc/testsuite/gcc.dg/gomp/target-3.c new file mode 100644 index 0000000..3e79212 --- /dev/null +++ b/gcc/testsuite/gcc.dg/gomp/target-3.c @@ -0,0 +1,16 @@ +/* { dg-do compile } */ +/* { dg-options "-fopenmp -fdump-tree-gimple" } */ + +struct S +{ + int a, b; +}; + +void foo (struct S *s) +{ + #pragma omp target map (alloc: s->a, s->b) + ; + #pragma omp target enter data map (alloc: s->a, s->b) +} + +/* { dg-final { scan-tree-dump-times "map\\(struct:\\*s \\\[len: 2\\\]\\) map\\(alloc:s->a \\\[len: \[0-9\]+\\\]\\) map\\(alloc:s->b \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */ diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c index fcc0796..a81ba40 100644 --- a/gcc/tree-pretty-print.c +++ b/gcc/tree-pretty-print.c @@ -858,6 +858,7 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) { case GOMP_MAP_ALLOC: case GOMP_MAP_POINTER: + case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION: pp_string (pp, "alloc"); break; case GOMP_MAP_IF_PRESENT: @@ -936,6 +937,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) case GOMP_MAP_ATTACH_DETACH: pp_string (pp, "attach_detach"); break; + case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION: + pp_string (pp, "attach_zero_length_array_section"); + break; default: gcc_unreachable (); } @@ -954,6 +958,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) case GOMP_MAP_ALWAYS_POINTER: pp_string (pp, " [pointer assign, bias: "); break; + case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION: + pp_string (pp, " [pointer assign, zero-length array section, bias: "); + break; case GOMP_MAP_TO_PSET: pp_string (pp, " [pointer set, len: "); break; @@ -961,6 +968,7 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) case GOMP_MAP_DETACH: case GOMP_MAP_FORCE_DETACH: case GOMP_MAP_ATTACH_DETACH: + case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION: pp_string (pp, " [bias: "); break; default: diff --git a/include/gomp-constants.h b/include/gomp-constants.h index 3e42d71..9e7db69 100644 --- a/include/gomp-constants.h +++ b/include/gomp-constants.h @@ -143,6 +143,11 @@ enum gomp_map_kind No refcount is bumped by this, and the store is done unconditionally. */ GOMP_MAP_ALWAYS_POINTER = (GOMP_MAP_FLAG_SPECIAL_2 | GOMP_MAP_FLAG_SPECIAL | 1), + /* Like GOMP_MAP_POINTER, but allow zero-length array section, i.e. set to + NULL if target is not mapped. */ + GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION + = (GOMP_MAP_FLAG_SPECIAL_2 + | GOMP_MAP_FLAG_SPECIAL | 2), /* Forced deallocation of zero length array section. */ GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION = (GOMP_MAP_FLAG_SPECIAL_2 @@ -163,6 +168,12 @@ enum gomp_map_kind GOMP_MAP_FORCE_DETACH = (GOMP_MAP_DEEP_COPY | GOMP_MAP_FLAG_FORCE | 1), + /* Like GOMP_MAP_ATTACH, but allow attaching to zero-length array sections + (i.e. set to NULL when array section is not mapped) Currently only used + by OpenMP. */ + GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION + = (GOMP_MAP_DEEP_COPY | 2), + /* Internal to GCC, not used in libgomp. */ /* Do not map, but pointer assign a pointer instead. */ GOMP_MAP_FIRSTPRIVATE_POINTER = (GOMP_MAP_LAST | 1), @@ -186,7 +197,8 @@ enum gomp_map_kind ((X) == GOMP_MAP_ALWAYS_POINTER) #define GOMP_MAP_POINTER_P(X) \ - ((X) == GOMP_MAP_POINTER) + ((X) == GOMP_MAP_POINTER \ + || (X) == GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION) #define GOMP_MAP_ALWAYS_TO_P(X) \ (((X) == GOMP_MAP_ALWAYS_TO) || ((X) == GOMP_MAP_ALWAYS_TOFROM)) diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h index 299cf42..832ca6e 100644 --- a/libgomp/libgomp.h +++ b/libgomp/libgomp.h @@ -1283,7 +1283,7 @@ extern uintptr_t gomp_map_val (struct target_mem_desc *, void **, size_t); extern void gomp_attach_pointer (struct gomp_device_descr *, struct goacc_asyncqueue *, splay_tree, splay_tree_key, uintptr_t, size_t, - struct gomp_coalesce_buf *); + struct gomp_coalesce_buf *, bool); extern void gomp_detach_pointer (struct gomp_device_descr *, struct goacc_asyncqueue *, splay_tree_key, uintptr_t, bool, struct gomp_coalesce_buf *); diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 5988db0..82d8dac 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -937,7 +937,7 @@ acc_attach_async (void **hostaddr, int async) } gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, n, (uintptr_t) hostaddr, - 0, NULL); + 0, NULL, false); gomp_mutex_unlock (&acc_dev->lock); } @@ -1141,7 +1141,7 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, if ((kinds[i] & 0xff) == GOMP_MAP_ATTACH) { gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, n, - (uintptr_t) h, s, NULL); + (uintptr_t) h, s, NULL, false); /* OpenACC 'attach'/'detach' doesn't affect structured/dynamic reference counts ('n->refcount', 'n->dynamic_refcount'). */ } @@ -1159,7 +1159,8 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, splay_tree_key m = lookup_host (acc_dev, hostaddrs[j], sizeof (void *)); gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, m, - (uintptr_t) hostaddrs[j], sizes[j], NULL); + (uintptr_t) hostaddrs[j], sizes[j], NULL, + false); } bool processed = false; diff --git a/libgomp/target.c b/libgomp/target.c index 5d3103a..6b02daf 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -543,7 +543,8 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep, struct gomp_coalesce_buf *cbuf, htab_t *refcount_set) { - assert (kind != GOMP_MAP_ATTACH); + assert (kind != GOMP_MAP_ATTACH + || kind != GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION); tgt_var->key = oldn; tgt_var->copy_from = GOMP_MAP_COPY_FROM_P (kind); @@ -616,7 +617,8 @@ get_implicit (bool short_mapkind, void *kinds, int idx) static void gomp_map_pointer (struct target_mem_desc *tgt, struct goacc_asyncqueue *aq, uintptr_t host_ptr, uintptr_t target_offset, uintptr_t bias, - struct gomp_coalesce_buf *cbuf) + struct gomp_coalesce_buf *cbuf, + bool allow_zero_length_array_sections) { struct gomp_device_descr *devicep = tgt->device_descr; struct splay_tree_s *mem_map = &devicep->mem_map; @@ -638,16 +640,24 @@ gomp_map_pointer (struct target_mem_desc *tgt, struct goacc_asyncqueue *aq, splay_tree_key n = gomp_map_lookup (mem_map, &cur_node); if (n == NULL) { - gomp_mutex_unlock (&devicep->lock); - gomp_fatal ("Pointer target of array section wasn't mapped"); - } - cur_node.host_start -= n->host_start; - cur_node.tgt_offset - = n->tgt->tgt_start + n->tgt_offset + cur_node.host_start; - /* At this point tgt_offset is target address of the - array section. Now subtract bias to get what we want - to initialize the pointer with. */ - cur_node.tgt_offset -= bias; + if (allow_zero_length_array_sections) + cur_node.tgt_offset = 0; + else + { + gomp_mutex_unlock (&devicep->lock); + gomp_fatal ("Pointer target of array section wasn't mapped"); + } + } + else + { + cur_node.host_start -= n->host_start; + cur_node.tgt_offset + = n->tgt->tgt_start + n->tgt_offset + cur_node.host_start; + /* At this point tgt_offset is target address of the + array section. Now subtract bias to get what we want + to initialize the pointer with. */ + cur_node.tgt_offset -= bias; + } gomp_copy_host2dev (devicep, aq, (void *) (tgt->tgt_start + target_offset), (void *) &cur_node.tgt_offset, sizeof (void *), true, cbuf); @@ -724,7 +734,8 @@ attribute_hidden void gomp_attach_pointer (struct gomp_device_descr *devicep, struct goacc_asyncqueue *aq, splay_tree mem_map, splay_tree_key n, uintptr_t attach_to, size_t bias, - struct gomp_coalesce_buf *cbufp) + struct gomp_coalesce_buf *cbufp, + bool allow_zero_length_array_sections) { struct splay_tree_key_s s; size_t size, idx; @@ -776,11 +787,19 @@ gomp_attach_pointer (struct gomp_device_descr *devicep, if (!tn) { - gomp_mutex_unlock (&devicep->lock); - gomp_fatal ("pointer target not mapped for attach"); + if (allow_zero_length_array_sections) + /* When allowing attachment to zero-length array sections, we + allow attaching to NULL pointers when the target region is not + mapped. */ + data = 0; + else + { + gomp_mutex_unlock (&devicep->lock); + gomp_fatal ("pointer target not mapped for attach"); + } } - - data = tn->tgt->tgt_start + tn->tgt_offset + target - tn->host_start; + else + data = tn->tgt->tgt_start + tn->tgt_offset + target - tn->host_start; gomp_debug (1, "%s: attaching host %p, target %p (struct base %p) to %p\n", @@ -1038,7 +1057,9 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, has_firstprivate = true; continue; } - else if ((kind & typemask) == GOMP_MAP_ATTACH) + else if ((kind & typemask) == GOMP_MAP_ATTACH + || ((kind & typemask) + == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION)) { tgt->list[i].key = NULL; has_firstprivate = true; @@ -1287,7 +1308,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, (uintptr_t) *(void **) hostaddrs[j], k->tgt_offset + ((uintptr_t) hostaddrs[j] - k->host_start), - sizes[j], cbufp); + sizes[j], cbufp, false); } } i = j - 1; @@ -1416,6 +1437,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, ++i; continue; case GOMP_MAP_ATTACH: + case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION: { cur_node.host_start = (uintptr_t) hostaddrs[i]; cur_node.host_end = cur_node.host_start + sizeof (void *); @@ -1432,9 +1454,12 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, structured/dynamic reference counts ('n->refcount', 'n->dynamic_refcount'). */ + bool zlas + = ((kind & typemask) + == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION); gomp_attach_pointer (devicep, aq, mem_map, n, (uintptr_t) hostaddrs[i], sizes[i], - cbufp); + cbufp, zlas); } else if ((pragma_kind & GOMP_MAP_VARS_OPENACC) != 0) { @@ -1545,9 +1570,12 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, false, cbufp); break; case GOMP_MAP_POINTER: - gomp_map_pointer (tgt, aq, - (uintptr_t) *(void **) k->host_start, - k->tgt_offset, sizes[i], cbufp); + case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION: + gomp_map_pointer + (tgt, aq, (uintptr_t) *(void **) k->host_start, + k->tgt_offset, sizes[i], cbufp, + ((kind & typemask) + == GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION)); break; case GOMP_MAP_TO_PSET: gomp_copy_host2dev (devicep, aq, @@ -1589,7 +1617,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, k->tgt_offset + ((uintptr_t) hostaddrs[j] - k->host_start), - sizes[j], cbufp); + sizes[j], cbufp, false); } } i = j - 1; diff --git a/libgomp/testsuite/libgomp.c++/target-23.C b/libgomp/testsuite/libgomp.c++/target-23.C new file mode 100644 index 0000000..d4f9ff3 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-23.C @@ -0,0 +1,34 @@ +extern "C" void abort (); + +struct S +{ + int *data; +}; + +int +main (void) +{ + #define SZ 10 + S *s = new S (); + s->data = new int[SZ]; + + for (int i = 0; i < SZ; i++) + s->data[i] = 0; + + #pragma omp target enter data map(to: s) + #pragma omp target enter data map(to: s->data[:SZ]) + #pragma omp target + { + for (int i = 0; i < SZ; i++) + s->data[i] = i; + } + #pragma omp target exit data map(from: s->data[:SZ]) + #pragma omp target exit data map(from: s) + + for (int i = 0; i < SZ; i++) + if (s->data[i] != i) + abort (); + + return 0; +} + diff --git a/libgomp/testsuite/libgomp.c++/target-lambda-1.C b/libgomp/testsuite/libgomp.c++/target-lambda-1.C new file mode 100644 index 0000000..06c6470 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-lambda-1.C @@ -0,0 +1,86 @@ +#include <cstdlib> +#include <cstring> + +template <typename L> +void +omp_target_loop (int begin, int end, L loop) +{ + #pragma omp target teams distribute parallel for + for (int i = begin; i < end; i++) + loop (i); +} + +struct S +{ + int a, len; + int *ptr; + + auto merge_data_func (int *iptr, int &b) + { + auto fn = [=](void) -> bool + { + bool mapped; + #pragma omp target map(from:mapped) + { + mapped = (ptr != NULL && iptr != NULL); + if (mapped) + { + for (int i = 0; i < len; i++) + ptr[i] += a + b + iptr[i]; + } + } + return mapped; + }; + return fn; + } +}; + +int x = 1; + +int main (void) +{ + const int N = 10; + int *data1 = new int[N]; + int *data2 = new int[N]; + memset (data1, 0xab, sizeof (int) * N); + memset (data1, 0xcd, sizeof (int) * N); + + int val = 1; + int &valref = val; + #pragma omp target enter data map(alloc: data1[:N], data2[:N]) + + omp_target_loop (0, N, [=](int i) { data1[i] = val; }); + omp_target_loop (0, N, [=](int i) { data2[i] = valref + 1; }); + + #pragma omp target update from(data1[:N], data2[:N]) + + for (int i = 0; i < N; i++) + { + if (data1[i] != 1) abort (); + if (data2[i] != 2) abort (); + } + + #pragma omp target exit data map(delete: data1[:N], data2[:N]) + + int b = 8; + S s = { 4, N, data1 }; + auto f = s.merge_data_func (data2, b); + + if (f ()) abort (); + + #pragma omp target enter data map(to: data1[:N]) + if (f ()) abort (); + + #pragma omp target enter data map(to: data2[:N]) + if (!f ()) abort (); + + #pragma omp target exit data map(from: data1[:N], data2[:N]) + + for (int i = 0; i < N; i++) + { + if (data1[i] != 0xf) abort (); + if (data2[i] != 2) abort (); + } + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/target-lambda-2.C b/libgomp/testsuite/libgomp.c++/target-lambda-2.C new file mode 100644 index 0000000..1d3561f --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-lambda-2.C @@ -0,0 +1,30 @@ +#include <cstdlib> + +#define N 10 +int main (void) +{ + int X, Y; + #pragma omp target map(from: X, Y) + { + int x = 0, y = 0; + + for (int i = 0; i < N; i++) + [&] (int v) { x += v; } (i); + + auto yinc = [&y] { y++; }; + for (int i = 0; i < N; i++) + yinc (); + + X = x; + Y = y; + } + + int Xs = 0; + for (int i = 0; i < N; i++) + Xs += i; + if (X != Xs) + abort (); + + if (Y != N) + abort (); +} diff --git a/libgomp/testsuite/libgomp.c++/target-this-1.C b/libgomp/testsuite/libgomp.c++/target-this-1.C new file mode 100644 index 0000000..a591ea4 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-this-1.C @@ -0,0 +1,29 @@ +extern "C" void abort (); + +struct S +{ + int a, b, c, d; + + int sum (void) + { + int val = 0; + val += a + b + this->c + this->d; + return val; + } + + int sum_offload (void) + { + int val = 0; + #pragma omp target map(val) + val += a + b + this->c + this->d; + return val; + } +}; + +int main (void) +{ + S s = { 1, 2, 3, 4 }; + if (s.sum () != s.sum_offload ()) + abort (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/target-this-2.C b/libgomp/testsuite/libgomp.c++/target-this-2.C new file mode 100644 index 0000000..8119be8 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-this-2.C @@ -0,0 +1,47 @@ + +// We use 'auto' without a function return type, so specify dialect here +// { dg-additional-options "-std=c++14" } + +extern "C" void abort (); + +struct T +{ + int x, y; + + auto sum_func (int n) + { + auto fn = [=](int m) -> int + { + int v; + v = (x + y) * n + m; + return v; + }; + return fn; + } + + auto sum_func_offload (int n) + { + auto fn = [=](int m) -> int + { + int v; + #pragma omp target map(from:v) + v = (x + y) * n + m; + return v; + }; + return fn; + } + +}; + +int main (void) +{ + T a = { 1, 2 }; + + auto s1 = a.sum_func (3); + auto s2 = a.sum_func_offload (3); + + if (s1 (1) != s2 (1)) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/target-this-3.C b/libgomp/testsuite/libgomp.c++/target-this-3.C new file mode 100644 index 0000000..e15f69a --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-this-3.C @@ -0,0 +1,99 @@ +#include <stdio.h> +#include <string.h> +extern "C" void abort (); + +struct S +{ + int * ptr; + int ptr_len; + + int *&refptr; + int refptr_len; + + bool set_ptr (int n) + { + bool mapped; + #pragma omp target map(from:mapped) + { + if (ptr != NULL) + for (int i = 0; i < ptr_len; i++) + ptr[i] = n; + mapped = (ptr != NULL); + } + return mapped; + } + + bool set_refptr (int n) + { + bool mapped; + #pragma omp target map(from:mapped) + { + if (refptr != NULL) + for (int i = 0; i < refptr_len; i++) + refptr[i] = n; + mapped = (refptr != NULL); + } + return mapped; + } +}; + +int main (void) +{ + #define N 10 + int *ptr1 = new int[N]; + int *ptr2 = new int[N]; + + memset (ptr1, 0, sizeof (int) * N); + memset (ptr2, 0, sizeof (int) * N); + + S s = { ptr1, N, ptr2, N }; + + bool mapped; + int val = 123; + + mapped = s.set_ptr (val); + if (mapped) + abort (); + if (s.ptr != ptr1) + abort (); + for (int i = 0; i < N; i++) + if (ptr1[i] != 0) + abort (); + + mapped = s.set_refptr (val); + if (mapped) + abort (); + if (s.refptr != ptr2) + abort (); + for (int i = 0; i < N; i++) + if (ptr2[i] != 0) + abort (); + + #pragma omp target data map(ptr1[:N]) + mapped = s.set_ptr (val); + + if (!mapped) + abort (); + if (s.set_refptr (0)) + abort (); + if (s.ptr != ptr1 || s.refptr != ptr2) + abort (); + for (int i = 0; i < N; i++) + if (ptr1[i] != val) + abort (); + + #pragma omp target data map(ptr2[:N]) + mapped = s.set_refptr (val); + + if (!mapped) + abort (); + if (s.set_ptr (0)) + abort (); + if (s.ptr != ptr1 || s.refptr != ptr2) + abort (); + for (int i = 0; i < N; i++) + if (ptr2[i] != val) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/target-this-4.C b/libgomp/testsuite/libgomp.c++/target-this-4.C new file mode 100644 index 0000000..9f53677 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-this-4.C @@ -0,0 +1,104 @@ + +// We use 'auto' without a function return type, so specify dialect here +// { dg-additional-options "-std=c++14" } +#include <cstdlib> +#include <cstring> + +struct T +{ + int *ptr; + int ptr_len; + + int *&refptr; + int refptr_len; + + auto set_ptr_func (int n) + { + auto fn = [=](void) -> bool + { + bool mapped; + #pragma omp target map(from:mapped) + { + if (ptr) + for (int i = 0; i < ptr_len; i++) + ptr[i] = n; + mapped = (ptr != NULL); + } + return mapped; + }; + return fn; + } + + auto set_refptr_func (int n) + { + auto fn = [=](void) -> bool + { + bool mapped; + #pragma omp target map(from:mapped) + { + if (refptr) + for (int i = 0; i < refptr_len; i++) + refptr[i] = n; + mapped = (refptr != NULL); + } + return mapped; + }; + return fn; + } +}; + +int main (void) +{ + #define N 10 + int *ptr1 = new int[N]; + int *ptr2 = new int[N]; + + memset (ptr1, 0, sizeof (int) * N); + memset (ptr2, 0, sizeof (int) * N); + + T a = { ptr1, N, ptr2, N }; + + auto p1 = a.set_ptr_func (1); + auto r2 = a.set_refptr_func (2); + + if (p1 ()) + abort (); + if (r2 ()) + abort (); + + if (a.ptr != ptr1) + abort (); + if (a.refptr != ptr2) + abort (); + + for (int i = 0; i < N; i++) + if (ptr1[i] != 0) + abort (); + + for (int i = 0; i < N; i++) + if (ptr2[i] != 0) + abort (); + + #pragma omp target data map(ptr1[:N], ptr2[:N]) + { + if (!p1 ()) + abort (); + if (!r2 ()) + abort (); + } + + if (a.ptr != ptr1) + abort (); + if (a.refptr != ptr2) + abort (); + + for (int i = 0; i < N; i++) + if (ptr1[i] != 1) + abort (); + + for (int i = 0; i < N; i++) + if (ptr2[i] != 2) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/target-this-5.C b/libgomp/testsuite/libgomp.c++/target-this-5.C new file mode 100644 index 0000000..e71c566 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-this-5.C @@ -0,0 +1,30 @@ +extern "C" void abort (); + +template<typename T> +struct S +{ + T a, b, c, d; + + T sum (void) + { + T val = 0; + val += a + b + this->c + this->d; + return val; + } + + T sum_offload (void) + { + T val = 0; + #pragma omp target map(val) + val += a + b + this->c + this->d; + return val; + } +}; + +int main (void) +{ + S<int> s = { 1, 2, 3, 4 }; + if (s.sum () != s.sum_offload ()) + abort (); + return 0; +} |