diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc index 8ca5954b6290..902bbfee9b97 100644 --- a/gcc/c-family/c-omp.cc +++ b/gcc/c-family/c-omp.cc @@ -3547,7 +3547,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 @@ -3709,6 +3711,80 @@ omp_expand_access_chain (tree *pc, tree expr, 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 * @@ -3723,7 +3799,7 @@ c_omp_address_inspector::expand_array_base (tree *pc, 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))); @@ -3734,6 +3810,7 @@ c_omp_address_inspector::expand_array_base (tree *pc, 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); @@ -3747,7 +3824,15 @@ c_omp_address_inspector::expand_array_base (tree *pc, return pc; } - switch (addr_tokens[i + 1]->u.access_kind) + 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[acc]->u.access_kind) { case ACCESS_DIRECT: if (decl_p && !target_p) @@ -4019,6 +4104,40 @@ c_omp_address_inspector::expand_array_base (tree *pc, } 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 NULL; @@ -4070,8 +4189,27 @@ c_omp_address_inspector::expand_component_selector (tree *pc, 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; - switch (addr_tokens[i + 1]->u.access_kind) + 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 + + 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: @@ -4081,7 +4219,7 @@ c_omp_address_inspector::expand_component_selector (tree *pc, { /* 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)); @@ -4090,7 +4228,7 @@ c_omp_address_inspector::expand_component_selector (tree *pc, 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; @@ -4101,15 +4239,15 @@ c_omp_address_inspector::expand_component_selector (tree *pc, 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, @@ -4126,12 +4264,12 @@ c_omp_address_inspector::expand_component_selector (tree *pc, 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, @@ -4146,10 +4284,10 @@ c_omp_address_inspector::expand_component_selector (tree *pc, 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); @@ -4164,13 +4302,38 @@ c_omp_address_inspector::expand_component_selector (tree *pc, /* ...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; + *idx = acc + 1; return NULL; } @@ -4188,8 +4351,7 @@ c_omp_address_inspector::expand_component_selector (tree *pc, pc = &OMP_CLAUSE_CHAIN (c); } - i += 2; - *idx = i; + *idx = acc + 1; if (chain_p && map_p) return omp_expand_access_chain (pc, expr, addr_tokens, idx, ort); diff --git a/gcc/c-family/c-pretty-print.cc b/gcc/c-family/c-pretty-print.cc index 1ce19f549887..1aac37b8e4a8 100644 --- a/gcc/c-family/c-pretty-print.cc +++ b/gcc/c-family/c-pretty-print.cc @@ -1644,6 +1644,11 @@ c_pretty_printer::postfix_expression (tree e) pp_colon (this); if (TREE_OPERAND (e, 2)) expression (TREE_OPERAND (e, 2)); + if (TREE_OPERAND (e, 3)) + { + pp_colon (this); + expression (TREE_OPERAND (e, 3)); + } pp_c_right_bracket (this); break; diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index 78a2f18a530b..d72cd10fbb8b 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -13633,7 +13633,7 @@ c_parser_postfix_expression_after_primary (c_parser *parser, len = c_parser_expression (parser).value; expr.value = build_omp_array_section (op_loc, expr.value, idx, - len); + len, NULL_TREE /* fixme */); } else expr.value = build_array_ref (op_loc, expr.value, idx); @@ -16339,11 +16339,11 @@ c_parser_oacc_wait_list (c_parser *parser, location_t clause_loc, tree list) struct omp_dim { - tree low_bound, length; + tree low_bound, length, stride; location_t loc; bool no_colon; - omp_dim (tree lb, tree len, location_t lo, bool nc) - : low_bound (lb), length (len), loc (lo), no_colon (nc) {} + omp_dim (tree lb, tree len, tree str, location_t lo, bool nc) + : low_bound (lb), length (len), stride (str), loc (lo), no_colon (nc) {} }; static tree @@ -16472,7 +16472,9 @@ c_parser_omp_variable_list (c_parser *parser, { tree low_bound = TREE_OPERAND (decl, 1); tree length = TREE_OPERAND (decl, 2); - dims.safe_push (omp_dim (low_bound, length, loc, false)); + tree stride = TREE_OPERAND (decl, 3); + dims.safe_push (omp_dim (low_bound, length, stride, loc, + false)); decl = TREE_OPERAND (decl, 0); } @@ -16488,21 +16490,22 @@ c_parser_omp_variable_list (c_parser *parser, else if (TREE_CODE (decl) == INDIRECT_REF) { dims.safe_push (omp_dim (integer_zero_node, - integer_one_node, loc, true)); + integer_one_node, NULL_TREE, loc, + true)); decl = TREE_OPERAND (decl, 0); } else /* ARRAY_REF. */ { tree index = TREE_OPERAND (decl, 1); - dims.safe_push (omp_dim (index, integer_one_node, loc, - true)); + dims.safe_push (omp_dim (index, integer_one_node, + NULL_TREE, loc, true)); decl = TREE_OPERAND (decl, 0); } } for (int i = dims.length () - 1; i >= 0; i--) decl = build_omp_array_section (loc, decl, dims[i].low_bound, - dims[i].length); + dims[i].length, dims[i].stride); } else if (TREE_CODE (decl) == INDIRECT_REF) { @@ -16511,7 +16514,7 @@ c_parser_omp_variable_list (c_parser *parser, STRIP_NOPS (decl); decl = build_omp_array_section (loc, decl, integer_zero_node, - integer_one_node); + integer_one_node, NULL_TREE); } else if (TREE_CODE (decl) == ARRAY_REF) { @@ -16520,7 +16523,8 @@ c_parser_omp_variable_list (c_parser *parser, decl = TREE_OPERAND (decl, 0); STRIP_NOPS (decl); - decl = build_omp_array_section (loc, decl, idx, integer_one_node); + decl = build_omp_array_section (loc, decl, idx, integer_one_node, + NULL_TREE); } else if (TREE_CODE (decl) == NON_LVALUE_EXPR || CONVERT_EXPR_P (decl)) @@ -16674,7 +16678,8 @@ c_parser_omp_variable_list (c_parser *parser, break; } - dims.safe_push (omp_dim (low_bound, length, loc, no_colon)); + dims.safe_push (omp_dim (low_bound, length, NULL_TREE, loc, + no_colon)); } if (t != error_mark_node) @@ -16698,7 +16703,8 @@ c_parser_omp_variable_list (c_parser *parser, for (unsigned i = 0; i < dims.length (); i++) t = build_omp_array_section (clause_loc, t, dims[i].low_bound, - dims[i].length); + dims[i].length, + dims[i].stride); } if ((kind == OMP_CLAUSE_DEPEND || kind == OMP_CLAUSE_AFFINITY) diff --git a/gcc/c/c-tree.h b/gcc/c/c-tree.h index 9fee3ba2791a..912b842ca13f 100644 --- a/gcc/c/c-tree.h +++ b/gcc/c/c-tree.h @@ -822,7 +822,7 @@ extern tree build_component_ref (location_t, tree, tree, location_t, location_t, bool = true); extern tree handle_counted_by_for_component_ref (location_t, tree); extern tree build_array_ref (location_t, tree, tree); -extern tree build_omp_array_section (location_t, tree, tree, tree); +extern tree build_omp_array_section (location_t, tree, tree, tree, tree); extern tree build_external_ref (location_t, tree, bool, tree *); extern void pop_maybe_used (bool); extern struct c_expr c_expr_sizeof_expr (location_t, struct c_expr); diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc index 731e41b448ad..61a5dc25cf57 100644 --- a/gcc/c/c-typeck.cc +++ b/gcc/c/c-typeck.cc @@ -2380,6 +2380,8 @@ mark_exp_read (tree exp) mark_exp_read (TREE_OPERAND (exp, 1)); if (TREE_OPERAND (exp, 2)) mark_exp_read (TREE_OPERAND (exp, 2)); + if (TREE_OPERAND (exp, 3)) + mark_exp_read (TREE_OPERAND (exp, 3)); break; default: break; @@ -3441,7 +3443,8 @@ build_array_ref (location_t loc, tree array, tree index) instead. */ tree -build_omp_array_section (location_t loc, tree array, tree index, tree length) +build_omp_array_section (location_t loc, tree array, tree index, tree length, + tree stride) { tree type = TREE_TYPE (array); gcc_assert (type); @@ -3478,7 +3481,8 @@ build_omp_array_section (location_t loc, tree array, tree index, tree length) sectype = c_build_array_type (eltype, idxtype); } - return build3_loc (loc, OMP_ARRAY_SECTION, sectype, array, index, length); + return build4_loc (loc, OMP_ARRAY_SECTION, sectype, array, index, length, + stride); } @@ -15181,7 +15185,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, bool &maybe_zero_len, unsigned int &first_non_one, bool &non_contiguous, enum c_omp_region_type ort) { - tree ret, low_bound, length, type; + tree ret, low_bound, length, stride, type; bool openacc = (ort & C_ORT_ACC) != 0; if (TREE_CODE (t) != OMP_ARRAY_SECTION) { @@ -15267,8 +15271,11 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, type = TREE_TYPE (ret); low_bound = TREE_OPERAND (t, 1); length = TREE_OPERAND (t, 2); + stride = TREE_OPERAND (t, 3); - if (low_bound == error_mark_node || length == error_mark_node) + if (low_bound == error_mark_node + || length == error_mark_node + || stride == error_mark_node) return error_mark_node; if (low_bound && !INTEGRAL_TYPE_P (TREE_TYPE (low_bound))) @@ -15285,6 +15292,13 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, length); return error_mark_node; } + if (stride && !INTEGRAL_TYPE_P (TREE_TYPE (stride))) + { + error_at (OMP_CLAUSE_LOCATION (c), + "stride %qE of array section does not have integral type", + stride); + return error_mark_node; + } if (low_bound && TREE_CODE (low_bound) == INTEGER_CST && TYPE_PRECISION (TREE_TYPE (low_bound)) @@ -15501,7 +15515,9 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, d = TREE_OPERAND (d, 0)) { tree d_length = TREE_OPERAND (d, 2); - if (d_length == NULL_TREE || !integer_onep (d_length)) + tree d_stride = TREE_OPERAND (d, 3); + if (d_length == NULL_TREE || !integer_onep (d_length) + || (d_stride && !integer_onep (d_stride))) { if (openacc && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP) { diff --git a/gcc/cp/cp-objcp-common.cc b/gcc/cp/cp-objcp-common.cc index 8336d0bb8f7c..f56541e8bc9b 100644 --- a/gcc/cp/cp-objcp-common.cc +++ b/gcc/cp/cp-objcp-common.cc @@ -685,6 +685,7 @@ cp_common_init_ts (void) MARK_TS_EXP (OFFSET_REF); MARK_TS_EXP (PSEUDO_DTOR_EXPR); MARK_TS_EXP (REINTERPRET_CAST_EXPR); + MARK_TS_EXP (OMP_ARRAYSHAPE_CAST_EXPR); MARK_TS_EXP (SCOPE_REF); MARK_TS_EXP (STATIC_CAST_EXPR); MARK_TS_EXP (STMT_EXPR); diff --git a/gcc/cp/cp-tree.def b/gcc/cp/cp-tree.def index bb5aaf983fee..bf79a525e072 100644 --- a/gcc/cp/cp-tree.def +++ b/gcc/cp/cp-tree.def @@ -257,6 +257,7 @@ DEFTREECODE (REINTERPRET_CAST_EXPR, "reinterpret_cast_expr", tcc_unary, 1) DEFTREECODE (CONST_CAST_EXPR, "const_cast_expr", tcc_unary, 1) DEFTREECODE (STATIC_CAST_EXPR, "static_cast_expr", tcc_unary, 1) DEFTREECODE (DYNAMIC_CAST_EXPR, "dynamic_cast_expr", tcc_unary, 1) +DEFTREECODE (OMP_ARRAYSHAPE_CAST_EXPR, "omp_arrayshape_cast_expr", tcc_unary, 1) DEFTREECODE (IMPLICIT_CONV_EXPR, "implicit_conv_expr", tcc_unary, 1) DEFTREECODE (DOTSTAR_EXPR, "dotstar_expr", tcc_expression, 2) DEFTREECODE (TYPEID_EXPR, "typeid_expr", tcc_expression, 1) diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h index 714357844fa5..5d7b3819ce1b 100644 --- a/gcc/cp/cp-tree.h +++ b/gcc/cp/cp-tree.h @@ -515,6 +515,7 @@ extern GTY(()) tree cp_global_trees[CPTI_MAX]; LOOKUP_FOUND_P (in RECORD_TYPE, UNION_TYPE, ENUMERAL_TYPE, NAMESPACE_DECL) FNDECL_MANIFESTLY_CONST_EVALUATED (in FUNCTION_DECL) TARGET_EXPR_INTERNAL_P (in TARGET_EXPR) + DECLTYPE_FOR_OMP_ARRAYSHAPE_CAST (in DECLTYPE_TYPE) 5: IDENTIFIER_VIRTUAL_P (in IDENTIFIER_NODE) FUNCTION_RVALUE_QUALIFIED (in FUNCTION_TYPE, METHOD_TYPE) CALL_EXPR_REVERSE_ARGS (in CALL_EXPR, AGGR_INIT_EXPR) @@ -5058,6 +5059,8 @@ get_vec_init_expr (tree t) TREE_LANG_FLAG_2 (DECLTYPE_TYPE_CHECK (NODE)) #define DECLTYPE_FOR_REF_CAPTURE(NODE) \ TREE_LANG_FLAG_3 (DECLTYPE_TYPE_CHECK (NODE)) +#define DECLTYPE_FOR_OMP_ARRAYSHAPE_CAST(NODE) \ + TREE_LANG_FLAG_4 (DECLTYPE_TYPE_CHECK (NODE)) /* Nonzero for VAR_DECL and FUNCTION_DECL node means that `extern' was specified in its declaration. This can also be set for an @@ -7229,6 +7232,8 @@ extern tree cxx_comdat_group (tree); extern bool cp_missing_noreturn_ok_p (tree); extern bool is_direct_enum_init (tree, tree); extern void initialize_artificial_var (tree, vec *); +extern tree cp_omp_create_arrayshape_type (location_t, tree, + vec *); extern tree check_var_type (tree, tree, location_t); extern tree reshape_init (tree, tree, tsubst_flags_t); extern tree next_aggregate_field (tree); @@ -7263,7 +7268,8 @@ extern void grokclassfn (tree, tree, enum overload_flags); extern tree grok_array_decl (location_t, tree, tree, vec **, tsubst_flags_t); -extern tree grok_omp_array_section (location_t, tree, tree, tree); +extern tree grok_omp_array_section (location_t, tree, tree, tree, + tree); extern tree delete_sanity (location_t, tree, tree, bool, int, tsubst_flags_t); extern tree check_classfn (tree, tree, tree); @@ -8159,6 +8165,8 @@ extern tree cp_build_vec_convert (tree, location_t, tree, tsubst_flags_t); extern tree cp_build_bit_cast (location_t, tree, tree, tsubst_flags_t); +extern tree cp_build_omp_arrayshape_cast (location_t, tree, tree, + tsubst_flags_t); extern void start_lambda_scope (tree decl); extern void finish_lambda_scope (void); extern void record_lambda_scope (tree lambda); @@ -8417,7 +8425,8 @@ inline tree build_x_binary_op (const op_location_t &loc, } extern tree build_x_array_ref (location_t, tree, tree, tsubst_flags_t); -extern tree build_omp_array_section (location_t, tree, tree, tree); +extern tree build_omp_array_section (location_t, tree, tree, tree, + tree); extern tree build_x_unary_op (location_t, enum tree_code, cp_expr, tree, tsubst_flags_t); diff --git a/gcc/cp/decl.cc b/gcc/cp/decl.cc index 4a0a1fa12918..132558f9568e 100644 --- a/gcc/cp/decl.cc +++ b/gcc/cp/decl.cc @@ -12415,6 +12415,81 @@ create_array_type_for_decl (tree name, tree type, tree size, location_t loc) return build_cplus_array_type (type, itype); } +/* Build an anonymous array of SIZE elements of ELTYPE. */ + +static tree +create_anon_array_type (location_t loc, tree eltype, tree size) +{ + if (eltype == error_mark_node || size == error_mark_node) + return error_mark_node; + + tree itype = compute_array_index_type_loc (loc, NULL_TREE, size, + tf_warning_or_error); + + if (type_uses_auto (eltype) + && variably_modified_type_p (itype, /*fn=*/NULL_TREE)) + { + sorry_at (loc, "variable-length array of %"); + return error_mark_node; + } + + return build_cplus_array_type (eltype, itype); +} + +/* Derive an array type for an OpenMP array-shaping operator given EXPR, which + is an expression that might have array refs or array sections postfixed + (e.g. "ptr[0:3:2][3:4]"), and OMP_SHAPE_DIMS, a vector of dimensions. */ + +tree +cp_omp_create_arrayshape_type (location_t loc, tree expr, + vec *omp_shape_dims) +{ + tree type, strip_sections = expr; + + while (TREE_CODE (strip_sections) == OMP_ARRAY_SECTION + || TREE_CODE (strip_sections) == ARRAY_REF) + strip_sections = TREE_OPERAND (strip_sections, 0); + + /* Determine the element type, either directly or by using + "decltype" of an expression representing an element to + figure it out later during template instantiation. */ + if (type_dependent_expression_p (expr)) + { + type = cxx_make_type (DECLTYPE_TYPE); + + DECLTYPE_TYPE_EXPR (type) + = build_min_nt_loc (loc, INDIRECT_REF, strip_sections); + DECLTYPE_FOR_OMP_ARRAYSHAPE_CAST (type) = true; + SET_TYPE_STRUCTURAL_EQUALITY (type); + } + else + { + type = TREE_TYPE (strip_sections); + + if (TREE_CODE (type) == REFERENCE_TYPE) + type = TREE_TYPE (type); + + if (TREE_CODE (type) != POINTER_TYPE) + { + error ("OpenMP array shaping operator with non-pointer argument"); + return error_mark_node; + } + + type = TREE_TYPE (type); + } + + int i; + cp_expr dim; + FOR_EACH_VEC_ELT_REVERSE (*omp_shape_dims, i, dim) + { + if (!type_dependent_expression_p (dim)) + dim = fold_convert (sizetype, dim); + type = create_anon_array_type (loc, type, dim); + } + + return type; +} + /* Returns the smallest location that is not UNKNOWN_LOCATION. */ static location_t diff --git a/gcc/cp/decl2.cc b/gcc/cp/decl2.cc index 8b6ec14c0d66..23b4bfe9c6c7 100644 --- a/gcc/cp/decl2.cc +++ b/gcc/cp/decl2.cc @@ -622,35 +622,39 @@ grok_array_decl (location_t loc, tree array_expr, tree index_exp, tree grok_omp_array_section (location_t loc, tree array_expr, tree index, - tree length) + tree length, tree stride) { tree orig_array_expr = array_expr; tree orig_index = index; tree orig_length = length; + tree orig_stride = stride; if (error_operand_p (array_expr) || error_operand_p (index) - || error_operand_p (length)) + || error_operand_p (length) + || error_operand_p (stride)) return error_mark_node; if (processing_template_decl && (type_dependent_expression_p (array_expr) || type_dependent_expression_p (index) - || type_dependent_expression_p (length))) - return build_min_nt_loc (loc, OMP_ARRAY_SECTION, array_expr, index, length); + || type_dependent_expression_p (length) + || type_dependent_expression_p (stride))) + return build_min_nt_loc (loc, OMP_ARRAY_SECTION, array_expr, index, length, stride); index = fold_non_dependent_expr (index); length = fold_non_dependent_expr (length); + stride = fold_non_dependent_expr (stride); /* NOTE: We can pass through invalidly-typed index/length fields here (e.g. if the user tries to use a floating-point index/length). This is diagnosed later in semantics.cc:handle_omp_array_sections_1. */ - tree expr = build_omp_array_section (loc, array_expr, index, length); + tree expr = build_omp_array_section (loc, array_expr, index, length, stride); if (processing_template_decl) expr = build_min_non_dep (OMP_ARRAY_SECTION, expr, orig_array_expr, - orig_index, orig_length); + orig_index, orig_length, orig_stride); return expr; } @@ -2834,6 +2838,7 @@ min_vis_expr_r (tree *tp, int */*walk_subtrees*/, void *data) case REINTERPRET_CAST_EXPR: case CONST_CAST_EXPR: case DYNAMIC_CAST_EXPR: + case OMP_ARRAYSHAPE_CAST_EXPR: case NEW_EXPR: case CONSTRUCTOR: case LAMBDA_EXPR: diff --git a/gcc/cp/error.cc b/gcc/cp/error.cc index 40c5bf18ecb1..f9af2becefd7 100644 --- a/gcc/cp/error.cc +++ b/gcc/cp/error.cc @@ -2619,6 +2619,11 @@ dump_expr (cxx_pretty_printer *pp, tree t, int flags) dump_expr (pp, TREE_OPERAND (t, 1), flags); pp_colon (pp); dump_expr (pp, TREE_OPERAND (t, 2), flags); + if (TREE_OPERAND (t, 3)) + { + pp_colon (pp); + dump_expr (pp, TREE_OPERAND (t, 3), flags); + } pp_cxx_right_bracket (pp); break; diff --git a/gcc/cp/mangle.cc b/gcc/cp/mangle.cc index 3d5e96b2c944..6d0ce22398fe 100644 --- a/gcc/cp/mangle.cc +++ b/gcc/cp/mangle.cc @@ -3944,6 +3944,7 @@ write_expression (tree expr) case REINTERPRET_CAST_EXPR: case STATIC_CAST_EXPR: case CONST_CAST_EXPR: + case OMP_ARRAYSHAPE_CAST_EXPR: write_type (TREE_TYPE (expr)); write_expression (TREE_OPERAND (expr, 0)); break; diff --git a/gcc/cp/operators.def b/gcc/cp/operators.def index 17601d31207a..e24714bb2574 100644 --- a/gcc/cp/operators.def +++ b/gcc/cp/operators.def @@ -134,6 +134,7 @@ DEF_OPERATOR (NULL, DYNAMIC_CAST_EXPR, "dc", OVL_OP_FLAG_UNARY) DEF_OPERATOR (NULL, REINTERPRET_CAST_EXPR, "rc", OVL_OP_FLAG_UNARY) DEF_OPERATOR (NULL, CONST_CAST_EXPR, "cc", OVL_OP_FLAG_UNARY) DEF_OPERATOR (NULL, STATIC_CAST_EXPR, "sc", OVL_OP_FLAG_UNARY) +DEF_OPERATOR (NULL, OMP_ARRAYSHAPE_CAST_EXPR, "oc", OVL_OP_FLAG_UNARY) DEF_OPERATOR (NULL, SCOPE_REF, "sr", OVL_OP_FLAG_NONE) DEF_OPERATOR (NULL, EXPR_PACK_EXPANSION, "sp", OVL_OP_FLAG_NONE) DEF_OPERATOR (NULL, UNARY_LEFT_FOLD_EXPR, "fl", OVL_OP_FLAG_NONE) diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index 23aeb941eca0..ed578831c395 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -4652,6 +4652,12 @@ cp_parser_new (cp_lexer *lexer) /* Disallow OpenMP array sections in expressions. */ parser->omp_array_section_p = false; + /* Disallow OpenMP array-shaping operator in expressions. */ + parser->omp_array_shaping_op_p = false; + + /* We don't have an OpenMP array shape here. */ + parser->omp_has_array_shape_p = false; + /* Not declaring an implicit function template. */ parser->auto_is_implicit_function_template_parm_p = false; parser->fully_implicit_function_template_p = false; @@ -5659,6 +5665,7 @@ cp_parser_statement_expr (cp_parser *parser) { cp_token_position start = cp_parser_start_tentative_firewall (parser); auto oas = make_temp_override (parser->omp_array_section_p, false); + auto aso = make_temp_override (parser->omp_array_shaping_op_p, false); /* Consume the '('. */ location_t start_loc = cp_lexer_peek_token (parser->lexer)->location; @@ -8722,7 +8729,7 @@ cp_parser_postfix_open_square_expression (cp_parser *parser, && cp_lexer_next_token_is (parser->lexer, CPP_COLON)) { cp_lexer_consume_token (parser->lexer); - tree length = NULL_TREE; + tree length = NULL_TREE, stride = NULL_TREE; if (cp_lexer_next_token_is_not (parser->lexer, CPP_CLOSE_SQUARE)) { if (cxx_dialect >= cxx23) @@ -8755,9 +8762,23 @@ cp_parser_postfix_open_square_expression (cp_parser *parser, /*warn_comma_p=*/warn_comma_subscript); } + if (cp_lexer_next_token_is (parser->lexer, CPP_COLON)) + { + cp_lexer_consume_token (parser->lexer); + /* We could check for C++-23 multidimensional/comma-separated + subscripts here, or not bother. */ + if (cp_lexer_next_token_is_not (parser->lexer, CPP_CLOSE_SQUARE)) + stride + = cp_parser_expression (parser, NULL, /*cast_p=*/false, + /*decltype_p=*/false, + /*warn_comma_p=*/warn_comma_subscript); + } + parser->colon_corrects_to_scope_p = saved_colon_corrects_to_scope_p; - if (index == error_mark_node || length == error_mark_node) + if (index == error_mark_node + || length == error_mark_node + || stride == error_mark_node) { cp_parser_skip_to_closing_square_bracket (parser); return error_mark_node; @@ -8766,7 +8787,7 @@ cp_parser_postfix_open_square_expression (cp_parser *parser, cp_parser_require (parser, CPP_CLOSE_SQUARE, RT_CLOSE_SQUARE); return grok_omp_array_section (input_location, postfix_expression, index, - length); + length, stride); } parser->colon_corrects_to_scope_p = saved_colon_corrects_to_scope_p; @@ -8774,11 +8795,23 @@ cp_parser_postfix_open_square_expression (cp_parser *parser, /* Look for the closing `]'. */ cp_parser_require (parser, CPP_CLOSE_SQUARE, RT_CLOSE_SQUARE); - /* Build the ARRAY_REF. */ - postfix_expression = grok_array_decl (loc, postfix_expression, - index, &expression_list, - tf_warning_or_error - | (decltype_p ? tf_decltype : 0)); + if (parser->omp_has_array_shape_p + && (expression_list.get () == NULL + || vec_safe_length (expression_list) == 1)) + /* If we have an array-shaping operator, we may not be able to represent + a well-formed ARRAY_REF here, because we are coercing the type of the + innermost array base and the original type may not be compatible. Use + the OMP_ARRAY_SECTION code instead. We also want to explicitly avoid + creating INDIRECT_REFs for pointer bases, because that can lead to + parsing ambiguities (see cp_parser_omp_var_list_no_open). */ + return grok_omp_array_section (loc, postfix_expression, index, + size_one_node, NULL_TREE); + else + /* Build the ARRAY_REF. */ + postfix_expression = grok_array_decl (loc, postfix_expression, + index, &expression_list, + tf_warning_or_error + | (decltype_p ? tf_decltype : 0)); /* When not doing offsetof, array references are not permitted in constant-expressions. */ @@ -9101,6 +9134,7 @@ cp_parser_parenthesized_expression_list (cp_parser* parser, vec *expression_list; bool saved_greater_than_is_operator_p; bool saved_omp_array_section_p; + bool saved_omp_array_shaping_op_p; /* Assume all the expressions will be constant. */ if (non_constant_p) @@ -9119,7 +9153,9 @@ cp_parser_parenthesized_expression_list (cp_parser* parser, parser->greater_than_is_operator_p = true; saved_omp_array_section_p = parser->omp_array_section_p; + saved_omp_array_shaping_op_p = parser->omp_array_shaping_op_p; parser->omp_array_section_p = false; + parser->omp_array_shaping_op_p = false; cp_expr expr (NULL_TREE); @@ -9203,6 +9239,7 @@ cp_parser_parenthesized_expression_list (cp_parser* parser, parser->greater_than_is_operator_p = saved_greater_than_is_operator_p; parser->omp_array_section_p = saved_omp_array_section_p; + parser->omp_array_shaping_op_p = saved_omp_array_shaping_op_p; return NULL; } } @@ -9210,6 +9247,7 @@ cp_parser_parenthesized_expression_list (cp_parser* parser, parser->greater_than_is_operator_p = saved_greater_than_is_operator_p; parser->omp_array_section_p = saved_omp_array_section_p; + parser->omp_array_shaping_op_p = saved_omp_array_shaping_op_p; return expression_list; } @@ -10505,6 +10543,8 @@ cp_parser_cast_expression (cp_parser *parser, bool address_p, bool cast_p, cp_expr expr (NULL_TREE); int cast_expression = 0; const char *saved_message; + auto_vec omp_shape_dims; + bool omp_array_shape_p = false; /* There's no way to know yet whether or not this is a cast. For example, `(int (3))' is a unary-expression, while `(int) @@ -10574,6 +10614,28 @@ cp_parser_cast_expression (cp_parser *parser, bool address_p, bool cast_p, that the call to cp_parser_error_occurred below returns true. */ if (!cast_expression) cp_parser_simulate_error (parser); + else if (parser->omp_array_shaping_op_p + && cp_lexer_next_token_is (parser->lexer, CPP_OPEN_SQUARE)) + { + auto oas = make_temp_override (parser->omp_array_section_p, false); + auto aso = make_temp_override (parser->omp_array_shaping_op_p, false); + + while (cp_lexer_next_token_is (parser->lexer, CPP_OPEN_SQUARE)) + { + cp_lexer_consume_token (parser->lexer); + cp_expr e = cp_parser_expression (parser); + if (e.get_value () == error_mark_node) + break; + omp_shape_dims.safe_push (e); + if (!cp_parser_require (parser, CPP_CLOSE_SQUARE, + RT_CLOSE_SQUARE)) + break; + } + cp_token *close_paren = parens.require_close (parser); + if (close_paren) + close_paren_loc = close_paren->location; + omp_array_shape_p = true; + } else { type_id_in_expr_sentinel s (parser); @@ -10593,6 +10655,10 @@ cp_parser_cast_expression (cp_parser *parser, bool address_p, bool cast_p, function returning T. */ if (!cp_parser_error_occurred (parser)) { + auto aso = make_temp_override (parser->omp_array_shaping_op_p, false); + auto as = make_temp_override (parser->omp_has_array_shape_p, + omp_array_shape_p); + /* Only commit if the cast-expression doesn't start with '++', '--', or '[' in C++11. */ if (cast_expression > 0) @@ -10606,6 +10672,24 @@ cp_parser_cast_expression (cp_parser *parser, bool address_p, bool cast_p, if (cp_parser_parse_definitely (parser)) { + if (omp_array_shape_p) + { + location_t cast_loc = make_location (open_paren_loc, + open_paren_loc, + expr.get_finish ()); + + type = cp_omp_create_arrayshape_type (cast_loc, expr, + &omp_shape_dims); + + /* Things rapidly get worse below if we carry on from here + with an erroneous type... */ + if (error_operand_p (type)) + return error_mark_node; + + return cp_build_omp_arrayshape_cast (cast_loc, type, expr, + tf_warning_or_error); + } + /* Warn about old-style casts, if so requested. */ if (warn_old_style_cast && !in_system_header_at (input_location) @@ -11776,6 +11860,7 @@ cp_parser_lambda_expression (cp_parser* parser) bool auto_is_implicit_function_template_parm_p = parser->auto_is_implicit_function_template_parm_p; bool saved_omp_array_section_p = parser->omp_array_section_p; + bool saved_omp_array_shaping_op_p = parser->omp_array_shaping_op_p; parser->num_template_parameter_lists = 0; parser->in_statement = 0; @@ -11785,6 +11870,7 @@ cp_parser_lambda_expression (cp_parser* parser) parser->implicit_template_scope = 0; parser->auto_is_implicit_function_template_parm_p = false; parser->omp_array_section_p = false; + parser->omp_array_shaping_op_p = false; /* Inside the lambda, outside unevaluated context do not apply. */ cp_evaluated ev; @@ -11839,6 +11925,7 @@ cp_parser_lambda_expression (cp_parser* parser) parser->auto_is_implicit_function_template_parm_p = auto_is_implicit_function_template_parm_p; parser->omp_array_section_p = saved_omp_array_section_p; + parser->omp_array_shaping_op_p = saved_omp_array_shaping_op_p; } /* This lambda shouldn't have any proxies left at this point. */ @@ -26892,6 +26979,7 @@ cp_parser_braced_list (cp_parser *parser, bool *non_constant_p /*=nullptr*/) tree initializer; location_t start_loc = cp_lexer_peek_token (parser->lexer)->location; auto oas = make_temp_override (parser->omp_array_section_p, false); + auto aso = make_temp_override (parser->omp_array_shaping_op_p, false); /* Consume the `{' token. */ matching_braces braces; @@ -38985,11 +39073,11 @@ check_no_duplicate_clause (tree clauses, enum omp_clause_code code, struct omp_dim { - tree low_bound, length; + tree low_bound, length, stride; location_t loc; bool no_colon; - omp_dim (tree lb, tree len, location_t lo, bool nc) - : low_bound (lb), length (len), loc (lo), no_colon (nc) {} + omp_dim (tree lb, tree len, tree str, location_t lo, bool nc) + : low_bound (lb), length (len), stride (str), loc (lo), no_colon (nc) {} }; static tree @@ -39022,10 +39110,22 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind, || kind == OMP_CLAUSE_FROM)) { auto s = make_temp_override (parser->omp_array_section_p, true); + auto o = make_temp_override (parser->omp_array_shaping_op_p, + (kind == OMP_CLAUSE_TO + || kind == OMP_CLAUSE_FROM)); + tree reshaped_to = NULL_TREE; token = cp_lexer_peek_token (parser->lexer); location_t loc = token->location; decl = cp_parser_assignment_expression (parser); + if ((TREE_CODE (decl) == VIEW_CONVERT_EXPR + && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) + || TREE_CODE (decl) == OMP_ARRAYSHAPE_CAST_EXPR) + { + reshaped_to = TREE_TYPE (decl); + decl = TREE_OPERAND (decl, 0); + } + /* This code rewrites a parsed expression containing various tree codes used to represent array accesses into a more uniform nest of OMP_ARRAY_SECTION nodes before it is processed by @@ -39036,49 +39136,159 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind, dims.truncate (0); if (TREE_CODE (decl) == OMP_ARRAY_SECTION) { + size_t sections = 0; + tree orig_decl = decl; + bool update_p = (kind == OMP_CLAUSE_TO + || kind == OMP_CLAUSE_FROM); + bool maybe_ptr_based_noncontig_update = false; + + while (update_p + && !reshaped_to + && (TREE_CODE (decl) == OMP_ARRAY_SECTION + || TREE_CODE (decl) == ARRAY_REF + || TREE_CODE (decl) == COMPOUND_EXPR)) + { + if (TREE_CODE (decl) == COMPOUND_EXPR) + decl = TREE_OPERAND (decl, 1); + else + { + if (TREE_CODE (decl) == OMP_ARRAY_SECTION) + maybe_ptr_based_noncontig_update = true; + decl = TREE_OPERAND (decl, 0); + sections++; + } + } + + decl = orig_decl; + while (TREE_CODE (decl) == OMP_ARRAY_SECTION) { tree low_bound = TREE_OPERAND (decl, 1); tree length = TREE_OPERAND (decl, 2); - dims.safe_push (omp_dim (low_bound, length, loc, false)); + tree stride = TREE_OPERAND (decl, 3); + dims.safe_push (omp_dim (low_bound, length, stride, loc, + false)); decl = TREE_OPERAND (decl, 0); + if (sections > 0) + sections--; } + /* The handling of INDIRECT_REF here in the presence of + array-shaping operations is a little tricky. We need to + avoid treating a pointer dereference as a unit-sized array + section when we have an array shaping operation, because we + don't want an indirection to consume one of the user's + requested array dimensions. E.g. if we have a + double-indirect pointer like: + + int **foopp; + #pragma omp target update from(([N][N]) (*foopp)[0:X][0:Y]) + + We don't want to interpret this as: + + foopp[0:1][0:X][0:Y] + + else the array shape [N][N] won't match. Also we can't match + the array sections right-to-left instead, else this: + + #pragma omp target update from(([N][N]) (*foopp)[0:X]) + + would not copy the dimensions: + + (*foopp)[0:X][0:N] + + as required. So, avoid descending through INDIRECT_REFs if + we have an array-shaping op. + + If we *don't* have an array-shaping op, but we have a + multiply-indirected pointer and an array section like this: + + int ***fooppp; + #pragma omp target update from((**fooppp)[0:X:S] + + also avoid descending through more indirections than we have + array sections, since the noncontiguous update processing code + won't understand them (and doesn't need to traverse them + anyway). */ + while (TREE_CODE (decl) == ARRAY_REF - || TREE_CODE (decl) == INDIRECT_REF + || (TREE_CODE (decl) == INDIRECT_REF + && !reshaped_to) || TREE_CODE (decl) == COMPOUND_EXPR) { if (REFERENCE_REF_P (decl)) break; + if (maybe_ptr_based_noncontig_update && sections == 0) + break; + if (TREE_CODE (decl) == COMPOUND_EXPR) { decl = TREE_OPERAND (decl, 1); STRIP_NOPS (decl); + continue; } - else if (TREE_CODE (decl) == INDIRECT_REF) + else if (TREE_CODE (decl) == INDIRECT_REF + && !reshaped_to) { dims.safe_push (omp_dim (integer_zero_node, - integer_one_node, loc, true)); + integer_one_node, NULL_TREE, loc, + true)); decl = TREE_OPERAND (decl, 0); } else /* ARRAY_REF. */ { tree index = TREE_OPERAND (decl, 1); - dims.safe_push (omp_dim (index, integer_one_node, loc, - true)); + dims.safe_push (omp_dim (index, integer_one_node, + NULL_TREE, loc, true)); decl = TREE_OPERAND (decl, 0); + if (sections > 0) + sections--; } } + if (reshaped_to) + { + unsigned reshaped_dims = 0; + + for (tree t = reshaped_to; + TREE_CODE (t) == ARRAY_TYPE; + t = TREE_TYPE (t)) + reshaped_dims++; + + if (dims.length () > reshaped_dims) + { + error_at (loc, "too many array section specifiers " + "for %qT", reshaped_to); + decl = error_mark_node; + } + else + { + /* We have a pointer DECL whose target should be + interpreted as an array with particular dimensions, + not "the pointer itself". So, add an indirection + here. */ + if (type_dependent_expression_p (decl)) + decl = build_min_nt_loc (loc, INDIRECT_REF, decl); + else + { + /* We're interested in the reference target. */ + decl = convert_from_reference (decl); + decl = cp_build_fold_indirect_ref (decl); + } + decl + = cp_build_omp_arrayshape_cast (loc, reshaped_to, decl, + tf_warning_or_error); + } + } /* Bare references have their own special handling, so remove the explicit dereference added by convert_from_reference. */ - if (REFERENCE_REF_P (decl)) + else if (REFERENCE_REF_P (decl)) decl = TREE_OPERAND (decl, 0); for (int i = dims.length () - 1; i >= 0; i--) decl = grok_omp_array_section (loc, decl, dims[i].low_bound, - dims[i].length); + dims[i].length, dims[i].stride); } else if (TREE_CODE (decl) == INDIRECT_REF) { @@ -39094,7 +39304,7 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind, if (!ref_p) decl = grok_omp_array_section (loc, decl, integer_zero_node, - integer_one_node); + integer_one_node, NULL_TREE); } else if (TREE_CODE (decl) == ARRAY_REF) { @@ -39103,7 +39313,16 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind, decl = TREE_OPERAND (decl, 0); STRIP_NOPS (decl); - decl = grok_omp_array_section (loc, decl, idx, integer_one_node); + decl = grok_omp_array_section (loc, decl, idx, integer_one_node, + NULL_TREE); + } + else if (reshaped_to) + { + /* We're copying the whole of a reshaped array, originally a + base pointer. Rewrite as an array section. */ + tree elems = array_type_nelts_total (reshaped_to); + decl = grok_omp_array_section (loc, decl, size_zero_node, elems, + NULL_TREE); } else if (TREE_CODE (decl) == NON_LVALUE_EXPR || CONVERT_EXPR_P (decl)) @@ -39268,7 +39487,8 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind, goto skip_comma; } - dims.safe_push (omp_dim (low_bound, length, loc, no_colon)); + dims.safe_push (omp_dim (low_bound, length, NULL_TREE, loc, + no_colon)); } if ((kind == OMP_CLAUSE_MAP @@ -39290,7 +39510,8 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind, for (unsigned i = 0; i < dims.length (); i++) decl = build_omp_array_section (input_location, decl, dims[i].low_bound, - dims[i].length); + dims[i].length, + dims[i].stride); break; default: break; @@ -39303,6 +39524,8 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind, && cp_parser_simulate_error (parser)) { depend_lvalue: + auto o = make_temp_override (parser->omp_array_shaping_op_p, + true); cp_parser_abort_tentative_parse (parser); decl = cp_parser_assignment_expression (parser, NULL, false, false); @@ -48996,8 +49219,38 @@ cp_parser_omp_target_update (cp_parser *parser, cp_token *pragma_tok, tree clauses = cp_parser_omp_all_clauses (parser, OMP_TARGET_UPDATE_CLAUSE_MASK, "#pragma omp target update", pragma_tok); - if (omp_find_clause (clauses, OMP_CLAUSE_TO) == NULL_TREE - && omp_find_clause (clauses, OMP_CLAUSE_FROM) == NULL_TREE) + bool to_clause = false, from_clause = false; + for (tree c = clauses; + c && !to_clause && !from_clause; + c = OMP_CLAUSE_CHAIN (c)) + { + switch (OMP_CLAUSE_CODE (c)) + { + case OMP_CLAUSE_TO: + to_clause = true; + break; + case OMP_CLAUSE_FROM: + from_clause = true; + break; + case OMP_CLAUSE_MAP: + switch (OMP_CLAUSE_MAP_KIND (c)) + { + case GOMP_MAP_TO_GRID: + to_clause = true; + break; + case GOMP_MAP_FROM_GRID: + from_clause = true; + break; + default: + ; + } + break; + default: + ; + } + } + + if (!to_clause && !from_clause) { error_at (pragma_tok->location, "%<#pragma omp target update%> must contain at least one " diff --git a/gcc/cp/parser.h b/gcc/cp/parser.h index f9ed80123c4c..7622eb46455c 100644 --- a/gcc/cp/parser.h +++ b/gcc/cp/parser.h @@ -419,6 +419,13 @@ struct GTY(()) cp_parser { /* TRUE if an OpenMP array section is allowed. */ bool omp_array_section_p; + /* TRUE if an OpenMP array-shaping operator is allowed. */ + bool omp_array_shaping_op_p; + + /* TRUE if we are parsing an expression with an OpenMP array-shaping + operator. */ + bool omp_has_array_shape_p; + /* Tracks the function's template parameter list when declaring a function using generic type parameters. This is either a new chain in the case of a fully implicit function template or an extension of the function's existing diff --git a/gcc/cp/pt.cc b/gcc/cp/pt.cc index a25c04b427e7..ea06740960af 100644 --- a/gcc/cp/pt.cc +++ b/gcc/cp/pt.cc @@ -17262,6 +17262,10 @@ tsubst (tree t, tree args, tsubst_flags_t complain, tree in_decl) member access. */ id = false; type = finish_decltype_type (type, id, complain); + + if (DECLTYPE_FOR_OMP_ARRAYSHAPE_CAST (t) + && TYPE_REF_P (type)) + type = TREE_TYPE (type); } return cp_build_qualified_type (type, cp_type_quals (t) @@ -17921,14 +17925,17 @@ tsubst_omp_clause_decl (tree decl, tree args, tsubst_flags_t complain, = tsubst_stmt (TREE_OPERAND (decl, 1), args, complain, in_decl); tree length = tsubst_stmt (TREE_OPERAND (decl, 2), args, complain, in_decl); + tree stride = tsubst_stmt (TREE_OPERAND (decl, 3), args, complain, + in_decl); tree base = tsubst_omp_clause_decl (TREE_OPERAND (decl, 0), args, complain, in_decl, NULL); if (TREE_OPERAND (decl, 0) == base && TREE_OPERAND (decl, 1) == low_bound - && TREE_OPERAND (decl, 2) == length) + && TREE_OPERAND (decl, 2) == length + && TREE_OPERAND (decl, 3) == stride) return decl; - return build3 (OMP_ARRAY_SECTION, TREE_TYPE (base), base, low_bound, - length); + return build4 (OMP_ARRAY_SECTION, TREE_TYPE (base), base, low_bound, + length, stride); } tree ret = tsubst_stmt (decl, args, complain, in_decl); /* Undo convert_from_reference tsubst_expr could have called. */ @@ -20827,6 +20834,14 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl) RETURN (cp_build_bit_cast (EXPR_LOCATION (t), type, op0, complain)); } + case OMP_ARRAYSHAPE_CAST_EXPR: + { + tree type = tsubst (TREE_TYPE (t), args, complain, in_decl); + tree op0 = RECUR (TREE_OPERAND (t, 0)); + RETURN (cp_build_omp_arrayshape_cast (EXPR_LOCATION (t), type, op0, + complain)); + } + case POSTDECREMENT_EXPR: case POSTINCREMENT_EXPR: op1 = tsubst_non_call_postfix_expression (TREE_OPERAND (t, 0), @@ -21012,7 +21027,7 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl) case OMP_ARRAY_SECTION: { tree op0 = RECUR (TREE_OPERAND (t, 0)); - tree op1 = NULL_TREE, op2 = NULL_TREE; + tree op1 = NULL_TREE, op2 = NULL_TREE, op3 = NULL_TREE; if (op0 == error_mark_node) RETURN (error_mark_node); if (TREE_OPERAND (t, 1)) @@ -21027,7 +21042,14 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl) if (op2 == error_mark_node) RETURN (error_mark_node); } - RETURN (build_omp_array_section (EXPR_LOCATION (t), op0, op1, op2)); + if (TREE_OPERAND (t, 3)) + { + op3 = RECUR (TREE_OPERAND (t, 3)); + if (op3 == error_mark_node) + RETURN (error_mark_node); + } + RETURN (build_omp_array_section (EXPR_LOCATION (t), op0, op1, op2, + op3)); } case OMP_DECLARE_MAPPER: diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc index 5ea63518c226..562736ede7a4 100644 --- a/gcc/cp/semantics.cc +++ b/gcc/cp/semantics.cc @@ -5914,9 +5914,10 @@ public: static tree handle_omp_array_sections_1 (tree c, tree t, vec &types, bool &maybe_zero_len, unsigned int &first_non_one, - bool &non_contiguous, enum c_omp_region_type ort) + bool &non_contiguous, enum c_omp_region_type ort, + int *discontiguous) { - tree ret, low_bound, length, type; + tree ret, low_bound, length, stride, type; bool openacc = (ort & C_ORT_ACC) != 0; if (TREE_CODE (t) != OMP_ARRAY_SECTION) { @@ -5979,18 +5980,25 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, TREE_OPERAND (t, 0) = omp_privatize_field (TREE_OPERAND (t, 0), false); ret = handle_omp_array_sections_1 (c, TREE_OPERAND (t, 0), types, maybe_zero_len, first_non_one, - non_contiguous, ort); + non_contiguous, ort, discontiguous); if (ret == error_mark_node || ret == NULL_TREE) return ret; - type = TREE_TYPE (ret); + if (TREE_CODE (ret) == OMP_ARRAY_SECTION) + type = TREE_TYPE (TREE_TYPE (TREE_OPERAND (ret, 0))); + else + type = TREE_TYPE (ret); low_bound = TREE_OPERAND (t, 1); length = TREE_OPERAND (t, 2); + stride = TREE_OPERAND (t, 3); if ((low_bound && type_dependent_expression_p (low_bound)) - || (length && type_dependent_expression_p (length))) + || (length && type_dependent_expression_p (length)) + || (stride && type_dependent_expression_p (stride))) return NULL_TREE; - if (low_bound == error_mark_node || length == error_mark_node) + if (low_bound == error_mark_node + || length == error_mark_node + || stride == error_mark_node) return error_mark_node; if (low_bound && !INTEGRAL_TYPE_P (TREE_TYPE (low_bound))) @@ -6007,15 +6015,26 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, length); return error_mark_node; } + if (stride && !INTEGRAL_TYPE_P (TREE_TYPE (stride))) + { + error_at (OMP_CLAUSE_LOCATION (c), + "stride %qE of array section does not have integral type", + stride); + return error_mark_node; + } if (low_bound) low_bound = mark_rvalue_use (low_bound); if (length) length = mark_rvalue_use (length); + if (stride) + stride = mark_rvalue_use (stride); /* We need to reduce to real constant-values for checks below. */ if (length) length = fold_simple (length); if (low_bound) low_bound = fold_simple (low_bound); + if (stride) + stride = fold_simple (stride); if (low_bound && TREE_CODE (low_bound) == INTEGER_CST && TYPE_PRECISION (TREE_TYPE (low_bound)) @@ -6026,9 +6045,15 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, && TYPE_PRECISION (TREE_TYPE (length)) > TYPE_PRECISION (sizetype)) length = fold_convert (sizetype, length); + if (stride + && TREE_CODE (stride) == INTEGER_CST + && TYPE_PRECISION (TREE_TYPE (stride)) + > TYPE_PRECISION (sizetype)) + stride = fold_convert (sizetype, stride); if (low_bound == NULL_TREE) low_bound = integer_zero_node; - + if (stride == NULL_TREE) + stride = size_one_node; if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)) @@ -6147,12 +6172,29 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, } if (length && TREE_CODE (length) == INTEGER_CST) { - if (tree_int_cst_lt (size, length)) + tree slength = length; + if (stride && TREE_CODE (stride) == INTEGER_CST) { - error_at (OMP_CLAUSE_LOCATION (c), - "length %qE above array section size " - "in %qs clause", length, - omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + slength = size_binop (MULT_EXPR, + fold_convert (sizetype, length), + fold_convert (sizetype, stride)); + slength = size_binop (MINUS_EXPR, + slength, + fold_convert (sizetype, stride)); + slength = size_binop (PLUS_EXPR, slength, size_one_node); + } + if (tree_int_cst_lt (size, slength)) + { + if (stride) + error_at (OMP_CLAUSE_LOCATION (c), + "length %qE with stride %qE above array " + "section size in %qs clause", length, stride, + omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + else + error_at (OMP_CLAUSE_LOCATION (c), + "length %qE above array section size " + "in %qs clause", length, + omp_clause_code_name[OMP_CLAUSE_CODE (c)]); return error_mark_node; } if (TREE_CODE (low_bound) == INTEGER_CST) @@ -6160,7 +6202,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, tree lbpluslen = size_binop (PLUS_EXPR, fold_convert (sizetype, low_bound), - fold_convert (sizetype, length)); + fold_convert (sizetype, slength)); if (TREE_CODE (lbpluslen) == INTEGER_CST && tree_int_cst_lt (size, lbpluslen)) { @@ -6230,7 +6272,10 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, d = TREE_OPERAND (d, 0)) { tree d_length = TREE_OPERAND (d, 2); - if (d_length == NULL_TREE || !integer_onep (d_length)) + tree d_stride = TREE_OPERAND (d, 3); + if (d_length == NULL_TREE + || !integer_onep (d_length) + || (d_stride && !integer_onep (d_stride))) { if (openacc && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP) { @@ -6250,10 +6295,15 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, return error_mark_node; } - error_at (OMP_CLAUSE_LOCATION (c), - "array section is not contiguous in %qs clause", - omp_clause_code_name[OMP_CLAUSE_CODE (c)]); - return error_mark_node; + if (discontiguous && *discontiguous) + *discontiguous = 2; + else + { + error_at (OMP_CLAUSE_LOCATION (c), + "array section is not contiguous in %qs clause", + omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + return error_mark_node; + } } } } @@ -6265,7 +6315,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, return error_mark_node; } if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND) - types.safe_push (TREE_TYPE (ret)); + types.safe_push (type); /* We will need to evaluate lb more than once. */ tree lb = cp_save_expr (low_bound); if (lb != low_bound) @@ -6284,15 +6334,45 @@ handle_omp_array_sections_1 (tree c, tree t, vec &types, OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IN_REDUCTION || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TASK_REDUCTION); - ret = grok_array_decl (OMP_CLAUSE_LOCATION (c), ret, low_bound, NULL, - tf_warning_or_error); + /* NOTE: Stride/length are discarded for affinity/depend here. */ + if (discontiguous + && *discontiguous + && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_AFFINITY + && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND) + ret = grok_omp_array_section (OMP_CLAUSE_LOCATION (c), ret, low_bound, + length, stride); + else + ret = grok_array_decl (OMP_CLAUSE_LOCATION (c), ret, low_bound, NULL, + tf_warning_or_error); return ret; } -/* Handle array sections for clause C. */ +/* We built a reference to an array section, but it turns out we only need a + set of ARRAY_REFs to the lower bound. Rewrite the node. */ + +static tree +omp_array_section_low_bound (location_t loc, tree node) +{ + if (TREE_CODE (node) == OMP_ARRAY_SECTION) + { + tree low_bound = TREE_OPERAND (node, 1); + tree ret + = omp_array_section_low_bound (loc, TREE_OPERAND (node, 0)); + return grok_array_decl (loc, ret, low_bound, NULL, tf_warning_or_error); + } + + return node; +} + +/* Handle array sections for clause C. On entry *DISCONTIGUOUS is 0 if array + section must be contiguous, 1 if it can be discontiguous, and in the latter + case it is set to 2 on exit if it is determined to be discontiguous during + the function's execution. PC points to the clause to be processed, and + *PNEXT to the last mapping node created, if passed as non-NULL. */ static bool -handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort) +handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort, + int *discontiguous) { tree c = *pc; bool maybe_zero_len = false; @@ -6308,7 +6388,7 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort) tp = &TREE_VALUE (*tp); tree first = handle_omp_array_sections_1 (c, *tp, types, maybe_zero_len, first_non_one, - non_contiguous, ort); + non_contiguous, ort, discontiguous); if (first == error_mark_node) return true; if (first == NULL_TREE) @@ -6350,6 +6430,8 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort) if (processing_template_decl && maybe_zero_len) return false; + bool higher_discontiguous = false; + for (i = num, t = OMP_CLAUSE_DECL (c); i > 0; t = TREE_OPERAND (t, 0)) { @@ -6357,6 +6439,7 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort) tree low_bound = TREE_OPERAND (t, 1); tree length = TREE_OPERAND (t, 2); + tree stride = TREE_OPERAND (t, 3); i--; if (low_bound @@ -6369,6 +6452,11 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort) && TYPE_PRECISION (TREE_TYPE (length)) > TYPE_PRECISION (sizetype)) length = fold_convert (sizetype, length); + if (stride + && TREE_CODE (stride) == INTEGER_CST + && TYPE_PRECISION (TREE_TYPE (stride)) + > TYPE_PRECISION (sizetype)) + stride = fold_convert (sizetype, stride); if (low_bound == NULL_TREE) low_bound = integer_zero_node; @@ -6378,10 +6466,50 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort) continue; } + if (stride == NULL_TREE) + stride = size_one_node; + if (discontiguous && *discontiguous) + { + /* This condition is similar to the error check below, but + whereas that checks for a definitely-discontiguous array + section in order to report an error (where such a section is + illegal), here we instead need to know if the array section + *may be* discontiguous so we can handle that case + appropriately (i.e. for rectangular "target update" + operations). */ + bool full_span = false; + if (length != NULL_TREE + && TREE_CODE (length) == INTEGER_CST + && TREE_CODE (types[i]) == ARRAY_TYPE + && TYPE_DOMAIN (types[i]) + && TYPE_MAX_VALUE (TYPE_DOMAIN (types[i])) + && TREE_CODE (TYPE_MAX_VALUE (TYPE_DOMAIN (types[i]))) + == INTEGER_CST) + { + tree size; + size = size_binop (PLUS_EXPR, + TYPE_MAX_VALUE (TYPE_DOMAIN (types[i])), + size_one_node); + if (tree_int_cst_equal (length, size)) + full_span = true; + } + + if (!integer_onep (stride) + || (higher_discontiguous + && (!integer_zerop (low_bound) + || !full_span))) + *discontiguous = 2; + + if (!integer_onep (stride) + || !integer_zerop (low_bound) + || !full_span) + higher_discontiguous = true; + } + if (!maybe_zero_len && i > first_non_one) { if (integer_nonzerop (low_bound)) - goto do_warn_noncontiguous; + goto is_noncontiguous; if (length != NULL_TREE && TREE_CODE (length) == INTEGER_CST && TYPE_DOMAIN (types[i]) @@ -6395,12 +6523,17 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort) size_one_node); if (!tree_int_cst_equal (length, size)) { - do_warn_noncontiguous: - error_at (OMP_CLAUSE_LOCATION (c), - "array section is not contiguous in %qs " - "clause", - omp_clause_code_name[OMP_CLAUSE_CODE (c)]); - return true; + is_noncontiguous: + if (discontiguous && *discontiguous) + *discontiguous = 2; + else + { + error_at (OMP_CLAUSE_LOCATION (c), + "array section is not contiguous in %qs " + "clause", + omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + return true; + } } } if (!processing_template_decl @@ -6517,6 +6650,9 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort) OMP_CLAUSE_DECL (c) = t; return false; } + if (discontiguous && *discontiguous != 2) + first = omp_array_section_low_bound (OMP_CLAUSE_LOCATION (c), + first); OMP_CLAUSE_DECL (c) = first; if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR) return false; @@ -6528,9 +6664,6 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort) 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) - return false; - if (TREE_CODE (first) == INDIRECT_REF) { /* Detect and skip adding extra nodes for pointer-to-member @@ -6557,6 +6690,10 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort) } } + if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP + && !(discontiguous && *discontiguous == 2)) + return false; + /* FIRST represents the first item of data that we are mapping. E.g. if we're mapping an array, FIRST might resemble "foo.bar.myarray[0]". */ @@ -6575,7 +6712,8 @@ handle_omp_array_sections (tree *pc, tree **pnext, enum c_omp_region_type ort) c = *pc; - if (ai.maybe_zero_length_array_section (c)) + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && ai.maybe_zero_length_array_section (c)) OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1; /* !!! If we're accessing a base decl via chained access @@ -7797,7 +7935,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) t = OMP_CLAUSE_DECL (c); if (TREE_CODE (t) == OMP_ARRAY_SECTION) { - if (handle_omp_array_sections (pc, NULL, ort)) + if (handle_omp_array_sections (pc, NULL, ort, NULL)) { remove = true; break; @@ -8959,7 +9097,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) if (TREE_CODE (t) == OMP_ARRAY_SECTION) { - if (handle_omp_array_sections (pc, NULL, ort)) + int discontiguous = 1; + if (handle_omp_array_sections (pc, NULL, ort, &discontiguous)) remove = true; else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND && (OMP_CLAUSE_DEPEND_KIND (c) @@ -9114,6 +9253,9 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) remove = true; break; } + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_DIM + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_STRIDE) + break; /* FALLTHRU */ case OMP_CLAUSE_TO: case OMP_CLAUSE_FROM: @@ -9128,8 +9270,11 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) grp_start_p = pc; grp_sentinel = OMP_CLAUSE_CHAIN (c); + int discontiguous + = (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO + || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM); tree *pnext = NULL; - if (handle_omp_array_sections (pc, &pnext, ort)) + if (handle_omp_array_sections (pc, &pnext, ort, &discontiguous)) remove = true; else { @@ -9725,7 +9870,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) t = OMP_CLAUSE_DECL (c); if (TREE_CODE (t) == OMP_ARRAY_SECTION) { - if (handle_omp_array_sections (pc, NULL, ort)) + if (handle_omp_array_sections (pc, NULL, ort, NULL)) remove = true; else { @@ -14387,6 +14532,45 @@ cp_build_bit_cast (location_t loc, tree type, tree arg, return ret; } +/* Build an OpenMP array-shape cast of ARG to TYPE. */ + +tree +cp_build_omp_arrayshape_cast (location_t loc, tree type, tree arg, + tsubst_flags_t complain) +{ + if (error_operand_p (type)) + return error_mark_node; + + if (!dependent_type_p (type) + && !complete_type_or_maybe_complain (type, NULL_TREE, complain)) + return error_mark_node; + + if (error_operand_p (arg)) + return error_mark_node; + + if (!type_dependent_expression_p (arg) && !dependent_type_p (type)) + { + if (!trivially_copyable_p (TREE_TYPE (arg))) + { + error_at (cp_expr_loc_or_loc (arg, loc), + "OpenMP array shape source type %qT " + "is not trivially copyable", TREE_TYPE (arg)); + return error_mark_node; + } + + /* A pointer to multi-dimensional array conversion isn't normally + allowed, but we force it here for array shape operators by creating + the node directly. We also want to avoid any overloaded conversions + the user might have defined, not that there are likely to be any. */ + return build1_loc (loc, VIEW_CONVERT_EXPR, type, arg); + } + + tree ret = build_min (OMP_ARRAYSHAPE_CAST_EXPR, type, arg); + SET_EXPR_LOCATION (ret, loc); + + return ret; +} + /* Diagnose invalid #pragma GCC unroll argument and adjust it if needed. */ diff --git a/gcc/cp/typeck.cc b/gcc/cp/typeck.cc index 88f8f34e8060..1ab06647723a 100644 --- a/gcc/cp/typeck.cc +++ b/gcc/cp/typeck.cc @@ -1636,6 +1636,9 @@ structural_comptypes (tree t1, tree t2, int strict) return false; if (DECLTYPE_FOR_LAMBDA_PROXY (t1) != DECLTYPE_FOR_LAMBDA_PROXY (t2)) return false; + if (DECLTYPE_FOR_OMP_ARRAYSHAPE_CAST (t1) + != DECLTYPE_FOR_OMP_ARRAYSHAPE_CAST (t2)) + return false; if (!cp_tree_equal (DECLTYPE_TYPE_EXPR (t1), DECLTYPE_TYPE_EXPR (t2))) return false; break; @@ -4865,12 +4868,12 @@ build_x_array_ref (location_t loc, tree arg1, tree arg2, tree build_omp_array_section (location_t loc, tree array_expr, tree index, - tree length) + tree length, tree stride) { if (TREE_CODE (array_expr) == TYPE_DECL || type_dependent_expression_p (array_expr)) - return build3_loc (loc, OMP_ARRAY_SECTION, NULL_TREE, array_expr, index, - length); + return build4_loc (loc, OMP_ARRAY_SECTION, NULL_TREE, array_expr, index, + length, stride); tree type = TREE_TYPE (array_expr); gcc_assert (type); @@ -4909,8 +4912,8 @@ build_omp_array_section (location_t loc, tree array_expr, tree index, sectype = build_array_type (eltype, idxtype); } - return build3_loc (loc, OMP_ARRAY_SECTION, sectype, array_expr, index, - length); + return build4_loc (loc, OMP_ARRAY_SECTION, sectype, array_expr, index, + length, stride); } /* Return whether OP is an expression of enum type cast to integer @@ -8299,6 +8302,9 @@ check_for_casting_away_constness (location_t loc, tree src_type, src_type, dest_type); return true; + case OMP_ARRAYSHAPE_CAST_EXPR: + return true; + default: gcc_unreachable(); } diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index cfd3295075b9..f52e36274627 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -10629,6 +10629,19 @@ omp_group_last (tree *start_p) grp_last_p = &OMP_CLAUSE_CHAIN (c); break; + case GOMP_MAP_TO_GRID: + case GOMP_MAP_FROM_GRID: + while (nc + && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_DIM + || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_STRIDE)) + { + grp_last_p = &OMP_CLAUSE_CHAIN (c); + c = nc; + nc = OMP_CLAUSE_CHAIN (c); + } + break; + case GOMP_MAP_STRUCT: case GOMP_MAP_STRUCT_UNORD: { @@ -10777,6 +10790,10 @@ omp_group_base (omp_mapping_group *grp, unsigned int *chained, internal_error ("unexpected mapping node"); return error_mark_node; + case GOMP_MAP_TO_GRID: + case GOMP_MAP_FROM_GRID: + return *grp->grp_start; + case GOMP_MAP_ATTACH: case GOMP_MAP_DETACH: node = OMP_CLAUSE_CHAIN (node); @@ -15819,7 +15836,9 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, } if (remove) break; - if (OMP_CLAUSE_SIZE (c) == NULL_TREE) + if (OMP_CLAUSE_SIZE (c) == NULL_TREE + && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_GRID_DIM + && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_GRID_STRIDE) { /* Sanity check: attach/detach map kinds use the size as a bias, and it's never right to use the decl size for such @@ -15917,6 +15936,20 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, is_gimple_lvalue, fb_lvalue) == GS_ERROR) remove = true; } + else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_DIM + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_STRIDE) + { + /* The OMP_CLAUSE_DECL for GRID_DIM/GRID_STRIDE isn't necessarily + an lvalue -- e.g. it might be a constant. So handle it + specially here. */ + if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p, NULL, + is_gimple_val, fb_rvalue) == GS_ERROR) + { + gimplify_omp_ctxp = ctx; + remove = true; + } + break; + } else if (!DECL_P (decl)) { if ((ctx->region_type & ORT_TARGET) != 0 @@ -16009,8 +16042,13 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, gimplify_omp_ctxp = ctx->outer_context; if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, - fb_lvalue) == GS_ERROR) - remove = true; + fb_lvalue | fb_mayfail) == GS_ERROR) + { + sorry_at (OMP_CLAUSE_LOCATION (c), + "unsupported map expression %qE", + OMP_CLAUSE_DECL (c)); + remove = true; + } gimplify_omp_ctxp = ctx; break; } diff --git a/gcc/omp-general.cc b/gcc/omp-general.cc index 9ccbf6b417fe..5818ac618e9c 100644 --- a/gcc/omp-general.cc +++ b/gcc/omp-general.cc @@ -4114,6 +4114,32 @@ omp_addr_token::omp_addr_token (token_type t, structure_base_kinds k, tree e) u.structure_base_kind = k; } +static bool +omp_parse_noncontiguous_array (tree *expr0) +{ + tree expr = *expr0; + bool noncontig = false; + + while (TREE_CODE (expr) == OMP_ARRAY_SECTION + || TREE_CODE (expr) == ARRAY_REF) + { + /* Contiguous arrays use ARRAY_REF. By the time we reach here, + OMP_ARRAY_SECTION is only used for noncontiguous arrays. */ + if (TREE_CODE (expr) == OMP_ARRAY_SECTION) + noncontig = true; + + expr = TREE_OPERAND (expr, 0); + } + + if (noncontig) + { + *expr0 = expr; + return true; + } + + return false; +} + static bool omp_parse_component_selector (tree *expr0) { @@ -4213,6 +4239,13 @@ omp_parse_access_method (tree *expr0, enum access_method_kinds *kind) if (omp_parse_ref (&expr)) *kind = ACCESS_REF; + else if (omp_parse_noncontiguous_array (&expr)) + { + if (omp_parse_ref (&expr)) + *kind = ACCESS_NONCONTIG_REF_TO_ARRAY; + else + *kind = ACCESS_NONCONTIG_ARRAY; + } else if (omp_parse_pointer (&expr, &has_offset)) { if (omp_parse_ref (&expr)) @@ -4284,6 +4317,14 @@ omp_parse_structure_base (vec &addr_tokens, return true; } + if (TREE_CODE (expr) == VIEW_CONVERT_EXPR + && TREE_CODE (TREE_TYPE (expr)) == ARRAY_TYPE) + { + *kind = BASE_DECL; + *expr0 = TREE_OPERAND (expr, 0); + return true; + } + *kind = BASE_ARBITRARY_EXPR; *expr0 = expr; return true; @@ -4432,6 +4473,12 @@ debug_omp_tokenized_addr (vec &addr_tokens, case ACCESS_INDEXED_REF_TO_ARRAY: fputs ("access_indexed_ref_to_array", stderr); break; + case ACCESS_NONCONTIG_ARRAY: + fputs ("access_noncontig_array", stderr); + break; + case ACCESS_NONCONTIG_REF_TO_ARRAY: + fputs ("access_noncontig_ref_to_array", stderr); + break; } break; case ARRAY_BASE: diff --git a/gcc/omp-general.h b/gcc/omp-general.h index 37e331b2cfd4..ee9ae59b13d8 100644 --- a/gcc/omp-general.h +++ b/gcc/omp-general.h @@ -356,7 +356,9 @@ enum access_method_kinds ACCESS_POINTER_OFFSET, ACCESS_REF_TO_POINTER_OFFSET, ACCESS_INDEXED_ARRAY, - ACCESS_INDEXED_REF_TO_ARRAY + ACCESS_INDEXED_REF_TO_ARRAY, + ACCESS_NONCONTIG_ARRAY, + ACCESS_NONCONTIG_REF_TO_ARRAY }; /* These are the kinds that a STRUCTURE_BASE or ARRAY_BASE (except diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index 6123e7a214a0..5a150ee18e17 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -1272,6 +1272,55 @@ fixup_child_record_type (omp_context *ctx) : build_reference_type (type), TYPE_QUAL_RESTRICT); } +/* Build record type for noncontiguous target update operations. Must be kept + in sync with libgomp/libgomp.h omp_noncontig_array_desc. */ + +static tree +omp_noncontig_descriptor_type (location_t loc) +{ + static tree cached = NULL_TREE; + + if (cached) + return cached; + + tree t = make_node (RECORD_TYPE); + + tree fields = build_decl (loc, FIELD_DECL, get_identifier ("__ndims"), + size_type_node); + + tree field = build_decl (loc, FIELD_DECL, get_identifier ("__elemsize"), + size_type_node); + TREE_CHAIN (field) = fields; + fields = field; + + tree ptr_size_type = build_pointer_type (size_type_node); + + field = build_decl (loc, FIELD_DECL, get_identifier ("__dim"), ptr_size_type); + TREE_CHAIN (field) = fields; + fields = field; + + field = build_decl (loc, FIELD_DECL, get_identifier ("__index"), + ptr_size_type); + TREE_CHAIN (field) = fields; + fields = field; + + field = build_decl (loc, FIELD_DECL, get_identifier ("__length"), + ptr_size_type); + TREE_CHAIN (field) = fields; + fields = field; + + field = build_decl (loc, FIELD_DECL, get_identifier ("__stride"), + ptr_size_type); + TREE_CHAIN (field) = fields; + fields = field; + + finish_builtin_struct (t, "__omp_noncontig_desc_type", fields, ptr_type_node); + + cached = t; + + return t; +} + /* Instantiate decls as necessary in CTX to satisfy the data sharing specified by CLAUSES. */ @@ -1862,8 +1911,74 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) install_var_local (array_decl, ctx); break; } + else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_GRID + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FROM_GRID)) + { + tree desc_type = omp_noncontig_descriptor_type (UNKNOWN_LOCATION); - if (DECL_P (decl)) + tree bare = decl; + if (TREE_CODE (bare) == VIEW_CONVERT_EXPR) + bare = TREE_OPERAND (bare, 0); + + const char *desc_name = ".omp_noncontig_desc"; + /* Try (but not too hard) to make a friendly name for the + descriptor. */ + if (DECL_P (bare)) + desc_name = ACONCAT ((".omp_nc_desc_", + IDENTIFIER_POINTER (DECL_NAME (bare)), + NULL)); + tree desc = create_tmp_var (desc_type, desc_name); + DECL_NAMELESS (desc) = 1; + TREE_ADDRESSABLE (desc) = 1; + + /* Adjust DECL so it refers to the first element of the array: + either by indirecting a pointer, or by selecting the zero'th + index of each dimension of an array. (We don't have a "bias" + as such for this type of noncontiguous update operation, just + the volume specified in the descriptor we build in + lower_omp_target.) */ + + if (TREE_CODE (TREE_TYPE (decl)) == POINTER_TYPE) + { + decl = build_fold_indirect_ref (decl); + OMP_CLAUSE_DECL (c) = decl; + } + + tree field + = build_decl (OMP_CLAUSE_LOCATION (c), FIELD_DECL, NULL_TREE, + ptr_type_node); + SET_DECL_ALIGN (field, TYPE_ALIGN (ptr_type_node)); + insert_field_into_struct (ctx->record_type, field); + splay_tree_insert (ctx->field_map, (splay_tree_key) c, + (splay_tree_value) field); + + tree dn = build_omp_clause (OMP_CLAUSE_LOCATION (c), + OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (dn, GOMP_MAP_TO_PSET); + OMP_CLAUSE_DECL (dn) = desc; + OMP_CLAUSE_SIZE (dn) = TYPE_SIZE_UNIT (desc_type); + + OMP_CLAUSE_CHAIN (dn) = OMP_CLAUSE_CHAIN (c); + OMP_CLAUSE_CHAIN (c) = dn; + + field = build_decl (OMP_CLAUSE_LOCATION (c), FIELD_DECL, + NULL_TREE, ptr_type_node); + SET_DECL_ALIGN (field, TYPE_ALIGN (ptr_type_node)); + insert_field_into_struct (ctx->record_type, field); + splay_tree_insert (ctx->field_map, (splay_tree_key) dn, + (splay_tree_value) field); + + c = dn; + tree nc; + + while ((nc = OMP_CLAUSE_CHAIN (c)) + && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_DIM + || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_STRIDE)) + c = nc; + } + else if (DECL_P (decl)) { if (DECL_SIZE (decl) && !poly_int_tree_p (DECL_SIZE (decl))) @@ -2103,6 +2218,11 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) && is_omp_target (ctx->stmt) && !is_gimple_omp_offloaded (ctx->stmt)) break; + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_GRID + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FROM_GRID + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_DIM + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_STRIDE) + break; if (DECL_P (decl)) { if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER @@ -13252,6 +13372,10 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) case GOMP_MAP_DETACH: case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION: case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION: + case GOMP_MAP_TO_GRID: + case GOMP_MAP_FROM_GRID: + case GOMP_MAP_GRID_DIM: + case GOMP_MAP_GRID_STRIDE: break; case GOMP_MAP_IF_PRESENT: case GOMP_MAP_FORCE_ALLOC: @@ -13279,6 +13403,20 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) gcc_unreachable (); } #endif + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_GRID + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FROM_GRID) + { + tree nc = OMP_CLAUSE_CHAIN (c); + gcc_assert (OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_TO_PSET); + c = nc; + while ((nc = OMP_CLAUSE_CHAIN (c)) + && (OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_DIM + || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_STRIDE)) + c = nc; + map_cnt += 2; + continue; + } /* FALLTHRU */ case OMP_CLAUSE_TO: case OMP_CLAUSE_FROM: @@ -13698,7 +13836,267 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) deep_map_offset_data, deep_map_offset, &ilist); } - if (!DECL_P (ovar)) + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_GRID + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FROM_GRID)) + { + tree decl = OMP_CLAUSE_DECL (c); + tree dn = OMP_CLAUSE_CHAIN (c); + gcc_assert (OMP_CLAUSE_CODE (dn) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (dn) == GOMP_MAP_TO_PSET); + tree desc = OMP_CLAUSE_DECL (dn); + + tree oc, elsize = OMP_CLAUSE_SIZE (c); + tree type = TREE_TYPE (decl); + int i, dims = 0; + auto_vec tdims; + bool pointer_based = false, handled_pointer_section = false; + tree arrsize = fold_convert (sizetype, elsize); + + /* Allow a single (maybe strided) array section if we have a + pointer base. */ + if (TREE_CODE (decl) == INDIRECT_REF + && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0))) + == POINTER_TYPE)) + { + pointer_based = true; + dims = 1; + } + else + for (tree itype = type; + TREE_CODE (itype) == ARRAY_TYPE; + itype = TREE_TYPE (itype)) + { + tdims.safe_push (itype); + dims++; + } + + unsigned tdim = 0; + + vec *vdim; + vec *vindex; + vec *vlen; + vec *vstride; + vec_alloc (vdim, dims); + vec_alloc (vindex, dims); + vec_alloc (vlen, dims); + vec_alloc (vstride, dims); + + tree size_arr_type + = build_array_type_nelts (size_type_node, dims); + + tree dim_tmp = create_tmp_var (size_arr_type, ".omp_dim"); + DECL_NAMELESS (dim_tmp) = 1; + TREE_ADDRESSABLE (dim_tmp) = 1; + TREE_STATIC (dim_tmp) = 1; + tree index_tmp = create_tmp_var (size_arr_type, ".omp_index"); + DECL_NAMELESS (index_tmp) = 1; + TREE_ADDRESSABLE (index_tmp) = 1; + TREE_STATIC (index_tmp) = 1; + tree len_tmp = create_tmp_var (size_arr_type, ".omp_len"); + DECL_NAMELESS (len_tmp) = 1; + TREE_ADDRESSABLE (len_tmp) = 1; + TREE_STATIC (len_tmp) = 1; + tree stride_tmp = create_tmp_var (size_arr_type, ".omp_stride"); + DECL_NAMELESS (stride_tmp) = 1; + TREE_ADDRESSABLE (stride_tmp) = 1; + TREE_STATIC (stride_tmp) = 1; + + oc = c; + c = dn; + + for (i = 0; i < dims; i++) + { + nc = OMP_CLAUSE_CHAIN (c); + tree dim = NULL_TREE, index = NULL_TREE, len = NULL_TREE, + stride = size_one_node; + + if (OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_GRID_DIM) + { + index = OMP_CLAUSE_DECL (nc); + len = OMP_CLAUSE_SIZE (nc); + + index = fold_convert (sizetype, index); + len = fold_convert (sizetype, len); + + tree nc2 = OMP_CLAUSE_CHAIN (nc); + if (nc2 + && OMP_CLAUSE_CODE (nc2) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (nc2) + == GOMP_MAP_GRID_STRIDE)) + { + stride = OMP_CLAUSE_DECL (nc2); + stride = fold_convert (sizetype, stride); + nc = nc2; + } + + if (tdim < tdims.length ()) + { + /* We have an array shape -- use that to find the + total size of the data on the target to look up + in libgomp. */ + tree dtype = TYPE_DOMAIN (tdims[tdim]); + tree minval = TYPE_MIN_VALUE (dtype); + tree maxval = TYPE_MAX_VALUE (dtype); + minval = fold_convert (sizetype, minval); + maxval = fold_convert (sizetype, maxval); + dim = size_binop (MINUS_EXPR, maxval, minval); + dim = size_binop (PLUS_EXPR, dim, + size_one_node); + arrsize = size_binop (MULT_EXPR, arrsize, dim); + } + else if (pointer_based && !handled_pointer_section) + { + /* Use the selected array section to determine the + size of the array. */ + tree tmp = size_binop (MULT_EXPR, len, stride); + tmp = size_binop (MINUS_EXPR, tmp, stride); + tmp = size_binop (PLUS_EXPR, tmp, size_one_node); + dim = size_binop (PLUS_EXPR, index, tmp); + arrsize = size_binop (MULT_EXPR, arrsize, dim); + handled_pointer_section = true; + } + else + { + if (pointer_based) + error_at (OMP_CLAUSE_LOCATION (c), + "too many array section specifiers " + "for pointer-based array"); + else + error_at (OMP_CLAUSE_LOCATION (c), + "too many array section specifiers " + "for array"); + dim = index = len = stride = error_mark_node; + } + tdim++; + + c = nc; + } + else + { + /* We have more array dimensions than array section + specifiers. Copy the whole span. */ + tree dtype = TYPE_DOMAIN (tdims[tdim]); + tree minval = TYPE_MIN_VALUE (dtype); + tree maxval = TYPE_MAX_VALUE (dtype); + minval = fold_convert (sizetype, minval); + maxval = fold_convert (sizetype, maxval); + dim = size_binop (MINUS_EXPR, maxval, minval); + dim = size_binop (PLUS_EXPR, dim, size_one_node); + len = dim; + index = size_zero_node; + } + + if (TREE_CODE (dim) != INTEGER_CST) + TREE_STATIC (dim_tmp) = 0; + + if (TREE_CODE (index) != INTEGER_CST) + TREE_STATIC (index_tmp) = 0; + + if (TREE_CODE (len) != INTEGER_CST) + TREE_STATIC (len_tmp) = 0; + + if (TREE_CODE (stride) != INTEGER_CST) + TREE_STATIC (stride_tmp) = 0; + + tree cidx = size_int (i); + CONSTRUCTOR_APPEND_ELT (vdim, cidx, dim); + CONSTRUCTOR_APPEND_ELT (vindex, cidx, index); + CONSTRUCTOR_APPEND_ELT (vlen, cidx, len); + CONSTRUCTOR_APPEND_ELT (vstride, cidx, stride); + } + + /* The size of the whole array -- to make sure we find any + part of the array via splay-tree lookup that might be + mapped on the target at runtime. */ + OMP_CLAUSE_SIZE (oc) = arrsize; + + tree cdim = build_constructor (size_arr_type, vdim); + tree cindex = build_constructor (size_arr_type, vindex); + tree clen = build_constructor (size_arr_type, vlen); + tree cstride = build_constructor (size_arr_type, vstride); + + if (TREE_STATIC (dim_tmp)) + DECL_INITIAL (dim_tmp) = cdim; + else + gimplify_assign (dim_tmp, cdim, &ilist); + + if (TREE_STATIC (index_tmp)) + DECL_INITIAL (index_tmp) = cindex; + else + gimplify_assign (index_tmp, cindex, &ilist); + + if (TREE_STATIC (len_tmp)) + DECL_INITIAL (len_tmp) = clen; + else + gimplify_assign (len_tmp, clen, &ilist); + + if (TREE_STATIC (stride_tmp)) + DECL_INITIAL (stride_tmp) = cstride; + else + gimplify_assign (stride_tmp, cstride, &ilist); + + tree desc_type = TREE_TYPE (desc); + + tree ndims_field = TYPE_FIELDS (desc_type); + tree elemsize_field = DECL_CHAIN (ndims_field); + tree dim_field = DECL_CHAIN (elemsize_field); + tree index_field = DECL_CHAIN (dim_field); + tree len_field = DECL_CHAIN (index_field); + tree stride_field = DECL_CHAIN (len_field); + + vec *v; + vec_alloc (v, 6); + + bool all_static = (TREE_STATIC (dim_tmp) + && TREE_STATIC (index_tmp) + && TREE_STATIC (len_tmp) + && TREE_STATIC (stride_tmp)); + + dim_tmp = build4 (ARRAY_REF, sizetype, dim_tmp, size_zero_node, + NULL_TREE, NULL_TREE); + dim_tmp = build_fold_addr_expr (dim_tmp); + + /* TODO: we could skip all-zeros index. */ + index_tmp = build4 (ARRAY_REF, sizetype, index_tmp, + size_zero_node, NULL_TREE, NULL_TREE); + index_tmp = build_fold_addr_expr (index_tmp); + + len_tmp = build4 (ARRAY_REF, sizetype, len_tmp, size_zero_node, + NULL_TREE, NULL_TREE); + len_tmp = build_fold_addr_expr (len_tmp); + + /* TODO: we could skip all-ones stride. */ + stride_tmp = build4 (ARRAY_REF, sizetype, stride_tmp, + size_zero_node, NULL_TREE, NULL_TREE); + stride_tmp = build_fold_addr_expr (stride_tmp); + + elsize = fold_convert (sizetype, elsize); + tree ndims = size_int (dims); + + CONSTRUCTOR_APPEND_ELT (v, ndims_field, ndims); + CONSTRUCTOR_APPEND_ELT (v, elemsize_field, elsize); + CONSTRUCTOR_APPEND_ELT (v, dim_field, dim_tmp); + CONSTRUCTOR_APPEND_ELT (v, index_field, index_tmp); + CONSTRUCTOR_APPEND_ELT (v, len_field, len_tmp); + CONSTRUCTOR_APPEND_ELT (v, stride_field, stride_tmp); + + tree desc_ctor = build_constructor (desc_type, v); + + if (all_static) + { + TREE_STATIC (desc) = 1; + DECL_INITIAL (desc) = desc_ctor; + } + else + gimplify_assign (desc, desc_ctor, &ilist); + + OMP_CLAUSE_CHAIN (dn) = OMP_CLAUSE_CHAIN (nc); + c = oc; + nc = c; + } + else if (!DECL_P (ovar)) { if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)) diff --git a/gcc/testsuite/g++.dg/gomp/array-shaping-1.C b/gcc/testsuite/g++.dg/gomp/array-shaping-1.C new file mode 100644 index 000000000000..8627aa7ffb35 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/array-shaping-1.C @@ -0,0 +1,22 @@ +// { dg-do compile } +// { dg-additional-options "-fdump-tree-original" } + +template +void foo () +{ + T *ptr; + E a = A, b = B, c = C, d = D; + + /* Dependent types for indices. */ +#pragma omp target update from(([a][b+1][c][d]) ptr[1:a-2][1:b][1:c-2][1:d-2]) +// { dg-final { scan-tree-dump {map\(from_grid:VIEW_CONVERT_EXPR.*\(\*ptr\) \[len: 1\]\) map\(grid_dim:1 \[len: [^\]]+\]\) map\(grid_dim:1 \[len: [^\]]+\]\) map\(grid_dim:1 \[len: [^\]]+\]\) map\(grid_dim:1 \[len: [^]]+\]\)} "original" } } +} + +int main() +{ + char *ptr; + + foo (); + + return 0; +} diff --git a/gcc/testsuite/g++.dg/gomp/array-shaping-2.C b/gcc/testsuite/g++.dg/gomp/array-shaping-2.C new file mode 100644 index 000000000000..861d66261a14 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/array-shaping-2.C @@ -0,0 +1,134 @@ +// { dg-do compile } +// { dg-additional-options "-fdump-tree-original" } + +template +struct St +{ + T ***ppptr; + T ***&rppptr; + + St(T ***p, T ***&rp) : ppptr(p), rppptr(rp) { } +}; + +template +void foo() +{ + A *ptr; + A **pptr = &ptr; + A ***ppptr = &pptr; + A ***&rppptr = ppptr; + +#pragma omp target update to(([10]) (**ppptr)[3:4:2]) +// { dg-final { scan-tree-dump {map\(to_grid:VIEW_CONVERT_EXPR\(\*\*\*ppptr\) \[len: [0-9]+\]\) map\(grid_dim:3 \[len: 4\]\) map\(grid_stride:2\)} "original" } } + +#pragma omp target update to(([10]) (**rppptr)[3:4:2]) +// { dg-final { scan-tree-dump {map\(to_grid:VIEW_CONVERT_EXPR\(\*\*\*\*rppptr\) \[len: [0-9]+\]\) map\(grid_dim:3 \[len: 4\]\) map\(grid_stride:2\)} "original" } } + +#pragma omp target update to((**ppptr)[3:4:2]) +// { dg-final { scan-tree-dump {map\(to_grid:\*\*ppptr \[len: [0-9]+\]\) map\(grid_dim:3 \[len: 4\]\) map\(grid_stride:2\)} "original" } } + +#pragma omp target update to((**rppptr)[3:4:2]) +// { dg-final { scan-tree-dump {map\(to_grid:\*\*\*rppptr \[len: [0-9]+\]\) map\(grid_dim:3 \[len: 4\]\) map\(grid_stride:2\)} "original" } } + + B *ptr2; + B **pptr2 = &ptr2; + B ***ppptr2 = &pptr2; + St *s = new St(ppptr2, ppptr2); + St **ps = &s; + St **&rps = ps; + +#pragma omp target update from(([10]) (**(*ps)->ppptr)[3:4:2]) +// { dg-final { scan-tree-dump {map\(from_grid:VIEW_CONVERT_EXPR\(\*\*\*\(\*ps\)->ppptr\) \[len: [0-9]+\]\) map\(grid_dim:3 \[len: 4\]\) map\(grid_stride:2\)} "original" } } + +#pragma omp target update from(([10]) (**(*rps)->rppptr)[3:4:2]) +// { dg-final { scan-tree-dump {map\(from_grid:VIEW_CONVERT_EXPR\(\*\*\*\*\(\*\*rps\)->rppptr\) \[len: [0-9]+\]\) map\(grid_dim:3 \[len: 4\]\) map\(grid_stride:2\)} "original" } } + +#pragma omp target update from((**(*ps)->ppptr)[3:4:2]) +// { dg-final { scan-tree-dump {map\(from_grid:\*\*\(\*ps\)->ppptr \[len: [0-9]+\]\) map\(grid_dim:3 \[len: 4\]\) map\(grid_stride:2\)} "original" } } + +#pragma omp target update from((**(*rps)->rppptr)[3:4:2]) +// { dg-final { scan-tree-dump {map\(from_grid:\*\*\*\(\*\*rps\)->rppptr \[len: [0-9]+\]\) map\(grid_dim:3 \[len: 4\]\) map\(grid_stride:2\)} "original" } } + + B arr[10][10]; + B (*parr)[10][10] = &arr; + B (**pparr2)[10][10] = &parr; + B (**&rpparr2)[10][10] = pparr2; + +#pragma omp target update from(**pparr2) +// { dg-final { scan-tree-dump {from\(\*NON_LVALUE_EXPR <\*pparr2> \[len: [0-9]+\]\)} "original" } } + +#pragma omp target update to((**pparr2)[1:5:2][3:4:2]) +// { dg-final { scan-tree-dump {map\(to_grid:\*\*pparr2 \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\) map\(grid_dim:3 \[len: 4\]\) map\(grid_stride:2\)} "original" } } + +#pragma omp target update from((**rpparr2)[1:5:2][3:4:2]) +// { dg-final { scan-tree-dump {map\(from_grid:\*\*\*rpparr2 \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\) map\(grid_dim:3 \[len: 4\]\) map\(grid_stride:2\)} "original" } } + + delete s; +} + +struct S +{ + short ***ppptr; + short ***&rppptr; + + S(short ***p, short ***&rp) : ppptr(p), rppptr(rp) { } +}; + +int main() +{ + char *ptr; + char **pptr = &ptr; + char ***ppptr = &pptr; + char ***&rppptr = ppptr; + +#pragma omp target update to(([10]) (**ppptr)[1:5:2]) +// { dg-final { scan-tree-dump {map\(to_grid:VIEW_CONVERT_EXPR\(\*\*\*ppptr\) \[len: 1\]\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\)} "original" } } + +#pragma omp target update to(([10]) (**rppptr)[1:5:2]) +// { dg-final { scan-tree-dump {map\(to_grid:VIEW_CONVERT_EXPR\(\*\*\*\*rppptr\) \[len: 1\]\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\)} "original" } } + +#pragma omp target update to((**ppptr)[1:5:2]) +// { dg-final { scan-tree-dump {map\(to_grid:\*\*ppptr \[len: 1\]\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\)} "original" } } + +#pragma omp target update to((**rppptr)[1:5:2]) +// { dg-final { scan-tree-dump {map\(to_grid:\*\*\*rppptr \[len: 1\]\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\)} "original" } } + + short *ptr2; + short **pptr2 = &ptr2; + short ***ppptr2 = &pptr2; + S *s = new S(ppptr2, ppptr2); + S **ps = &s; + S **&rps = ps; + +#pragma omp target update from(([10]) (**(*ps)->ppptr)[1:5:2]) +// { dg-final { scan-tree-dump {map\(from_grid:VIEW_CONVERT_EXPR\(\*\*\*\(\*ps\)->ppptr\) \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\)} "original" } } + +#pragma omp target update from(([10]) (**(*rps)->rppptr)[1:5:2]) +// { dg-final { scan-tree-dump {map\(from_grid:VIEW_CONVERT_EXPR\(\*\*\*\*\(\*\*rps\)->rppptr\) \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\)} "original" } } + +#pragma omp target update from((**(*ps)->ppptr)[1:5:2]) +// { dg-final { scan-tree-dump {map\(from_grid:\*\*\(\*ps\)->ppptr \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\)} "original" } } + +#pragma omp target update from((**(*rps)->rppptr)[1:5:2]) +// { dg-final { scan-tree-dump {map\(from_grid:\*\*\*\(\*\*rps\)->rppptr \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\)} "original" } } + + delete s; + + short arr[10][10]; + short (*parr)[10][10] = &arr; + short (**pparr)[10][10] = &parr; + short (**&rpparr)[10][10] = pparr; + +#pragma omp target update from(**pparr) +// { dg-final { scan-tree-dump {from\(\*NON_LVALUE_EXPR <\*pparr> \[len: [0-9]+\]\)} "original" } } + +#pragma omp target update to((**pparr)[1:5:2][1:5:2]) +// { dg-final { scan-tree-dump {map\(to_grid:\*\*pparr \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\)} "original" } } + +#pragma omp target update from((**rpparr)[1:5:2][1:5:2]) +// { dg-final { scan-tree-dump {map\(from_grid:\*\*\*rpparr \[len: [0-9]+\]\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\) map\(grid_dim:1 \[len: 5\]\) map\(grid_stride:2\)} "original" } } + + foo (); + + return 0; +} diff --git a/gcc/testsuite/g++.dg/gomp/bad-array-shaping-1.C b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-1.C new file mode 100644 index 000000000000..1f4e68bc065a --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-1.C @@ -0,0 +1,47 @@ +// { dg-do compile } + +#include +#include + +template +void foo (T *w) +{ + memset (w, 0, sizeof (T) * 100); + +#pragma omp target enter data map(to: w[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + w[j * 10 + i] = i + j * 3; + +#pragma omp target update to(([C][D]) w[3:2][1:8][0:5]) +// { dg-error "too many array section specifiers for" "" { target *-*-* } .-1 } +// { dg-error "'#pragma omp target update' must contain at least one 'from' or 'to' clauses" "" { target *-*-* } .-2 } + +#pragma omp target exit data map(from: w[:100]) +} + +int main() +{ + float *arr = new float[100]; + + memset (arr, 0, sizeof (float) * 100); + +#pragma omp target enter data map(to: arr[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + arr[j * 10 + i] = i + j * 3; + +#pragma omp target update to(([10][10]) arr[3:2][1:8][0:5]) +// { dg-error "too many array section specifiers for" "" { target *-*-* } .-1 } +// { dg-error "'#pragma omp target update' must contain at least one 'from' or 'to' clauses" "" { target *-*-* } .-2 } + +#pragma omp target exit data map(from: arr[:100]) + + foo (arr); + + delete[] arr; + + return 0; +} diff --git a/gcc/testsuite/g++.dg/gomp/bad-array-shaping-2.C b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-2.C new file mode 100644 index 000000000000..d32092925464 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-2.C @@ -0,0 +1,52 @@ +// { dg-do compile } + +#include +#include + +template +void foo (T *w) +{ + /* This isn't allowed. We get a cascade of errors because it looks a bit + like lambda-definition syntax */ +#pragma omp target enter data map(to: ([C][D]) w[:100]) + // { dg-error {capture of non-variable 'C'} "" { target *-*-* } .-1 } + // { dg-error {expected '\{' before '\[' token} "" { target *-*-* } .-2 } + // { dg-warning {lambda expressions only available with} "" { target c++98_only } .-3 } + // { dg-error {expected '\)' before 'w'} "" { target *-*-* } .-4 } + // { dg-error {does not have pointer or array type} "" { target *-*-* } .-5 } + +#pragma omp target exit data map(from: ([C][D]) w[:100]) + // { dg-error {capture of non-variable 'C'} "" { target *-*-* } .-1 } + // { dg-error {expected '\{' before '\[' token} "" { target *-*-* } .-2 } + // { dg-warning {lambda expressions only available with} "" { target c++98_only } .-3 } + // { dg-error {expected '\)' before 'w'} "" { target *-*-* } .-4 } + // { dg-error {does not have pointer or array type} "" { target *-*-* } .-5 } +} + +int main() +{ + float *arr = new float[100]; + + /* This isn't allowed (as above). */ +#pragma omp target enter data map(to: ([10][10]) arr[:100]) + // { dg-error {expected identifier before numeric constant} "" { target *-*-* } .-1 } + // { dg-error {expected '\{' before '\[' token} "" { target *-*-* } .-2 } + // { dg-warning {lambda expressions only available with} "" { target c++98_only } .-3 } + // { dg-error {expected '\)' before 'arr'} "" { target *-*-* } .-4 } + // { dg-error {no match for 'operator\[\]' in} "" { target *-*-* } .-5 } + // { dg-error {'#pragma omp target enter data' must contain at least one 'map' clause} "" { target *-*-*} .-6 } + +#pragma omp target exit data map(from: ([10][10]) arr[:100]) + // { dg-error {expected identifier before numeric constant} "" { target *-*-* } .-1 } + // { dg-error {expected '\{' before '\[' token} "" { target *-*-* } .-2 } + // { dg-warning {lambda expressions only available with} "" { target c++98_only } .-3 } + // { dg-error {no match for 'operator\[\]' in} "" { target *-*-* } .-4 } + // { dg-error {expected '\)' before 'arr'} "" { target *-*-* } .-5 } + // { dg-error {'#pragma omp target exit data' must contain at least one 'map' clause} "" { target *-*-* } .-6 } + + foo (arr); + + delete[] arr; + + return 0; +} diff --git a/gcc/testsuite/g++.dg/gomp/bad-array-shaping-3.C b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-3.C new file mode 100644 index 000000000000..90d0a5a80c52 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-3.C @@ -0,0 +1,53 @@ +// { dg-do compile } + +#include +#include + +template +void foo (T *w) +{ + memset (w, 0, sizeof (T) * 100); + int c = 50; + +#pragma omp target enter data map(to: w[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + w[j * 10 + i] = i + j * 3; + + /* This starts out looking like an array-shape cast. Make sure it's still + parsed as a lambda. */ +#pragma omp target update to(([c] (T *v) -> T { return v[c]; } (w))) + // { dg-message {sorry, unimplemented: unsupported map expression} "" { target *-*-* } .-1 } + // { dg-warning {lambda expressions only available with} "" { target c++98_only } .-2 } + +#pragma omp target exit data map(from: w[:100]) +} + +int main() +{ + float *arr = new float[100]; + int c = 50; + + memset (arr, 0, sizeof (float) * 100); + +#pragma omp target enter data map(to: arr[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + arr[j * 10 + i] = i + j * 3; + + /* As above. */ +#pragma omp target update to(([c] (float *v) -> float { return v[c]; } (arr))) + // { dg-message {sorry, unimplemented: unsupported map expression} "" { target *-*-* } .-1 } + // { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-2 } + // { dg-warning {lambda expressions only available with} "" { target c++98_only } .-3 } + +#pragma omp target exit data map(from: arr[:100]) + + foo (arr); + + delete[] arr; + + return 0; +} diff --git a/gcc/testsuite/g++.dg/gomp/bad-array-shaping-4.C b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-4.C new file mode 100644 index 000000000000..4518f03e9a0c --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-4.C @@ -0,0 +1,60 @@ +// { dg-do compile } + +#include +#include + +template +extern T* baz(T*); + +template +void foo (T *w) +{ + memset (w, 0, sizeof (T) * 100); + int c = 50; + +#pragma omp target enter data map(to: w[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + w[j * 10 + i] = i + j * 3; + + /* No array-shaping inside a function call. */ +#pragma omp target update to(baz(([10][10]) w)) + // { dg-error {expected identifier before numeric constant} "" { target *-*-* } .-1 } + // { dg-error {expected '\{' before '\[' token} "" { target *-*-* } .-2 } + // { dg-warning {lambda expressions only available with} "" { target c++98_only } .-3 } + // { dg-error {expected '\)' before 'w'} "" { target *-*-* } .-4 } + // { dg-error {no match for 'operator\[\]' in} "" { target *-*-* } .-5 } + +#pragma omp target exit data map(from: w[:100]) +} + +int main() +{ + float *arr = new float[100]; + int c = 50; + + memset (arr, 0, sizeof (float) * 100); + +#pragma omp target enter data map(to: arr[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + arr[j * 10 + i] = i + j * 3; + + /* As above. */ +#pragma omp target update to(baz(([10][10]) arr)) + // { dg-error {expected identifier before numeric constant} "" { target *-*-* } .-1 } + // { dg-error {expected '\{' before '\[' token} "" { target *-*-* } .-2 } + // { dg-warning {lambda expressions only available with} "" { target c++98_only } .-3 } + // { dg-error {no match for 'operator\[\]' in} "" { target *-*-* } .-4 } + // { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-5 } + +#pragma omp target exit data map(from: arr[:100]) + + foo (arr); + + delete[] arr; + + return 0; +} diff --git a/gcc/testsuite/g++.dg/gomp/bad-array-shaping-5.C b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-5.C new file mode 100644 index 000000000000..25edb9d1d9d3 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-5.C @@ -0,0 +1,55 @@ +// { dg-do compile } +// { dg-additional-options "-std=c++14" } + +#include +#include + +template +void foo (T *w) +{ + memset (w, 0, sizeof (T) * 100); + int c = 50; + +#pragma omp target enter data map(to: w[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + w[j * 10 + i] = i + j * 3; + + /* No array-shaping inside a lambda body. */ +#pragma omp target update to([&](const int d) -> auto& { return ([d][d]) w; } (10)) +// { dg-error {expected '\{' before '\[' token} "" { target *-*-* } .-1 } +// { dg-error {expected ';' before 'w'} "" { target *-*-* } .-2 } +// { dg-error {no match for 'operator\[\]' in} "" { target *-*-* } .-3 } + +#pragma omp target exit data map(from: w[:100]) +} + +int main() +{ + float *arr = new float[100]; + int c = 50; + + memset (arr, 0, sizeof (float) * 100); + +#pragma omp target enter data map(to: arr[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + arr[j * 10 + i] = i + j * 3; + + /* As above. */ +#pragma omp target update to([&](const int d) -> auto& { return ([d][d]) arr; } (10)) +// { dg-error {expected '\{' before '\[' token} "" { target *-*-* } .-1 } +// { dg-error {no match for 'operator\[\]' in} "" { target *-*-* } .-2 } +// { dg-error {expected ';' before 'arr'} "" { target *-*-* } .-3 } +// { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-4 } + +#pragma omp target exit data map(from: arr[:100]) + + foo (arr); + + delete[] arr; + + return 0; +} diff --git a/gcc/testsuite/g++.dg/gomp/bad-array-shaping-6.C b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-6.C new file mode 100644 index 000000000000..e796eaa39a3d --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-6.C @@ -0,0 +1,59 @@ +// { dg-do compile } + +#include +#include + +template +void foo (T *w) +{ + memset (w, 0, sizeof (T) * 100); + +#pragma omp target enter data map(to: w[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + w[j * 10 + i] = i + j * 3; + + /* No array-shaping inside a statement expression. */ +#pragma omp target update to( ({ int d = 10; ([d][d]) w; )} ) +// { dg-error {expected '\{' before '\[' token} "" { target *-*-* } .-1 } +// { dg-warning {lambda expressions only available with} "" { target c++98_only } .-2 } +// { dg-error {no match for 'operator\[\]'} "" { target *-*-* } .-3 } +// { dg-error {expected ';' before 'w'} "" { target *-*-* } .-4 } +// { dg-error {expected primary-expression before '\)' token} "" { target *-*-* } .-5 } +// { dg-error {expected '\)' before end of line} "" { target *-*-* } .-6 } +// { dg-message {sorry, unimplemented: unsupported map expression} "" { target *-*-* } .-7 } + +#pragma omp target exit data map(from: w[:100]) +} + +int main() +{ + float *arr = new float[100]; + + memset (arr, 0, sizeof (float) * 100); + +#pragma omp target enter data map(to: arr[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + arr[j * 10 + i] = i + j * 3; + + /* As above. */ +#pragma omp target update to( ({ int d = 10; ([d][d]) arr; )} ) +// { dg-error {expected '\{' before '\[' token} "" { target *-*-* } .-1 } +// { dg-warning {lambda expressions only available with} "" { target c++98_only } .-2 } +// { dg-error {no match for 'operator\[\]'} "" { target *-*-* } .-3 } +// { dg-error {expected primary-expression before '\)' token} "" { target *-*-* } .-4 } +// { dg-error {expected '\)' before end of line} "" { target *-*-* } .-5 } +// { dg-message {sorry, unimplemented: unsupported map expression} "" { target *-*-* } .-6 } +// { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-7 } + +#pragma omp target exit data map(from: arr[:100]) + + foo (arr); + + delete[] arr; + + return 0; +} diff --git a/gcc/testsuite/g++.dg/gomp/bad-array-shaping-7.C b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-7.C new file mode 100644 index 000000000000..c4b5d78b7f2a --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-7.C @@ -0,0 +1,44 @@ +// { dg-do compile } +// { dg-additional-options "-std=c++11" } + +#include + +template +struct St { + T *pp; +}; + +template +void foo (T *w) +{ + alignas (St) unsigned char buf[sizeof (St)]; + T *sub1; + + /* No array shaping op in brace initialiser (nonsensical anyway, but make + sure it doesn't parse). */ +#pragma omp target update to( new (buf) St { ([10][10]) sub1 } ) +// { dg-error {expected identifier before numeric constant} "" { target *-*-* } .-1 } +// { dg-error {expected '\{' before '\[' token} "" { target *-*-* } .-2 } +// { dg-error {expected '\}' before 'sub1'} "" { target *-*-* } .-3 } +} + +struct S { + int *pp; +}; + +int main() +{ + alignas (S) unsigned char buf[sizeof (S)]; + int *sub1; + + // As above. +#pragma omp target update to( new (buf) S { ([10][10]) sub1 } ) +// { dg-error {expected identifier before numeric constant} "" { target *-*-* } .-1 } +// { dg-error {expected '\{' before '\[' token} "" { target *-*-* } .-2 } +// { dg-error {expected '\}' before 'sub1'} "" { target *-*-* } .-3 } +// { dg-error {no match for 'operator\[\]'} "" { target *-*-* } .-4 } +// { dg-error {could not convert} "" { target *-*-* } .-5 } +// { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-6 } + + return 0; +} diff --git a/gcc/testsuite/g++.dg/gomp/bad-array-shaping-8.C b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-8.C new file mode 100644 index 000000000000..02d7de6088e0 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/bad-array-shaping-8.C @@ -0,0 +1,50 @@ +// { dg-do compile } + +template +void foo () +{ + T *ptr; + +#pragma omp target update to(([5][6][7]) ptr[0:4][0:7][0:7]) +// { dg-error {length '7' with stride '1' above array section size in 'to' clause} "" { target *-*-* } .-1 } + +#pragma omp target update to(([5][6][7]) ptr[1:5][0:6][0:7]) +// { dg-error {high bound '6' above array section size in 'to' clause} "" { target *-*-* } .-1 } + + // This one's OK... +#pragma omp target update from(([100]) ptr[3:33:3]) + + // But this is one element out of bounds. +#pragma omp target update from(([100]) ptr[4:33:3]) +// { dg-error {high bound '101' above array section size in 'from' clause} "" { target *-*-* } .-1 } + +#pragma omp target update to(([10][10]) ptr[0:9:-1][0:9]) +// { dg-error {length '9' with stride '-1' above array section size in 'to' clause} "" { target *-*-* } .-1 } +} + +int main() +{ + char *ptr; + +#pragma omp target update to(([5][6][7]) ptr[0:4][0:7][0:7]) +// { dg-error {length '7' with stride '1' above array section size in 'to' clause} "" { target *-*-* } .-1 } +// { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-2 } + +#pragma omp target update to(([5][6][7]) ptr[1:5][0:6][0:7]) +// { dg-error {high bound '6' above array section size in 'to' clause} "" { target *-*-* } .-1 } +// { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-2 } + +#pragma omp target update from(([100]) ptr[3:33:3]) + +#pragma omp target update from(([100]) ptr[4:33:3]) +// { dg-error {high bound '101' above array section size in 'from' clause} "" { target *-*-* } .-1 } +// { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-2 } + +#pragma omp target update to(([10][10]) ptr[0:9:-1][0:9]) +// { dg-error {length '9' with stride '-1' above array section size in 'to' clause} "" { target *-*-* } .-1 } +// { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-2 } + + foo (); + + return 0; +} diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc index 03e13e30b0a1..2826da7c2efe 100644 --- a/gcc/tree-pretty-print.cc +++ b/gcc/tree-pretty-print.cc @@ -1169,6 +1169,18 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) case GOMP_MAP_NONCONTIG_ARRAY_FORCE_PRESENT: pp_string (pp, "force_present,noncontig_array"); break; + case GOMP_MAP_TO_GRID: + pp_string (pp, "to_grid"); + break; + case GOMP_MAP_FROM_GRID: + pp_string (pp, "from_grid"); + break; + case GOMP_MAP_GRID_DIM: + pp_string (pp, "grid_dim"); + break; + case GOMP_MAP_GRID_STRIDE: + pp_string (pp, "grid_stride"); + break; case GOMP_MAP_UNSET: pp_string (pp, "unset"); break; @@ -2973,6 +2985,11 @@ dump_generic_node (pretty_printer *pp, tree node, int spc, dump_flags_t flags, dump_generic_node (pp, TREE_OPERAND (node, 1), spc, flags, false); pp_colon (pp); dump_generic_node (pp, TREE_OPERAND (node, 2), spc, flags, false); + if (TREE_OPERAND (node, 3)) + { + pp_colon (pp); + dump_generic_node (pp, TREE_OPERAND (node, 3), spc, flags, false); + } pp_right_bracket (pp); break; diff --git a/gcc/tree.def b/gcc/tree.def index 005f22f556a6..2253a679b749 100644 --- a/gcc/tree.def +++ b/gcc/tree.def @@ -1419,7 +1419,7 @@ DEFTREECODE (OMP_ATOMIC_CAPTURE_NEW, "omp_atomic_capture_new", tcc_statement, 2) DEFTREECODE (OMP_CLAUSE, "omp_clause", tcc_exceptional, 0) /* An OpenMP array section. */ -DEFTREECODE (OMP_ARRAY_SECTION, "omp_array_section", tcc_expression, 3) +DEFTREECODE (OMP_ARRAY_SECTION, "omp_array_section", tcc_expression, 4) /* OpenMP variant construct selector, used only in the middle end in the expansions of variant constructs that can't be resolved until the diff --git a/include/gomp-constants.h b/include/gomp-constants.h index 0b68e7774181..0bcfd2cf43c0 100644 --- a/include/gomp-constants.h +++ b/include/gomp-constants.h @@ -220,6 +220,9 @@ enum gomp_map_kind GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION = (GOMP_MAP_DEEP_COPY | 2), + GOMP_MAP_TO_GRID = (GOMP_MAP_DEEP_COPY | 4), + GOMP_MAP_FROM_GRID = (GOMP_MAP_DEEP_COPY | 5), + /* Internal to GCC, not used in libgomp. */ /* Do not map, but pointer assign a pointer instead. */ GOMP_MAP_FIRSTPRIVATE_POINTER = (GOMP_MAP_LAST | 1), @@ -243,7 +246,9 @@ enum gomp_map_kind GOMP_MAP_POP_MAPPER_NAME = (GOMP_MAP_LAST | 10), /* Used to hold a TREE_LIST of grouped nodes in an 'omp declare mapper' definition (only for Fortran at present). */ - GOMP_MAP_MAPPING_GROUP = (GOMP_MAP_LAST | 11) + GOMP_MAP_MAPPING_GROUP = (GOMP_MAP_LAST | 11), + GOMP_MAP_GRID_DIM = (GOMP_MAP_LAST | 12), + GOMP_MAP_GRID_STRIDE = (GOMP_MAP_LAST | 13) }; #define GOMP_MAP_COPY_TO_P(X) \ diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h index 3639f0014dc6..faf0c84c4d61 100644 --- a/libgomp/libgomp.h +++ b/libgomp/libgomp.h @@ -1333,6 +1333,20 @@ struct target_mem_desc { }; +/* A rectangular section of an array, for noncontiguous target update + operations. Must be kept in sync with + omp-low.cc:omp_noncontig_descriptor_type. */ + +typedef struct { + size_t ndims; + size_t elemsize; + size_t *dim; + size_t *index; + size_t *length; + size_t *stride; +} omp_noncontig_array_desc; + + typedef struct acc_dispatch_t { /* Execute. */ diff --git a/libgomp/target.c b/libgomp/target.c index 43f33c94c463..ca68faec3b8a 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -2319,6 +2319,14 @@ goacc_unmap_vars (struct target_mem_desc *tgt, bool do_copyfrom, gomp_unmap_vars_internal (tgt, do_copyfrom, NULL, aq); } +static int +omp_target_memcpy_rect_worker (void *, const void *, size_t, int, + const size_t *, const size_t *, const size_t *, + const size_t *, const size_t *, const size_t *, + struct gomp_device_descr *, + struct gomp_device_descr *, size_t *tmp_size, + void **tmp); + static void gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs, size_t *sizes, void *kinds, bool short_mapkind) @@ -2341,90 +2349,131 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs, } for (i = 0; i < mapnum; i++) - if (sizes[i]) - { - cur_node.host_start = (uintptr_t) hostaddrs[i]; - cur_node.host_end = cur_node.host_start + sizes[i]; - splay_tree_key n = splay_tree_lookup (&devicep->mem_map, &cur_node); - if (n) - { - int kind = get_kind (short_mapkind, kinds, i); - if (n->host_start > cur_node.host_start - || n->host_end < cur_node.host_end) - { - gomp_mutex_unlock (&devicep->lock); - gomp_fatal ("Trying to update [%p..%p) object when " - "only [%p..%p) is mapped", - (void *) cur_node.host_start, - (void *) cur_node.host_end, - (void *) n->host_start, - (void *) n->host_end); - } + { + int kind = get_kind (short_mapkind, kinds, i); + if ((kind & typemask) == GOMP_MAP_TO_GRID + || (kind & typemask) == GOMP_MAP_FROM_GRID) + { + omp_noncontig_array_desc *desc + = (omp_noncontig_array_desc *) hostaddrs[i + 1]; + cur_node.host_start = (uintptr_t) hostaddrs[i]; + cur_node.host_end = cur_node.host_start + sizes[i]; + assert (sizes[i + 1] == sizeof (omp_noncontig_array_desc)); + splay_tree_key n = splay_tree_lookup (&devicep->mem_map, &cur_node); + if (n) + { + if (n->aux && n->aux->attach_count) + { + gomp_mutex_unlock (&devicep->lock); + gomp_error ("noncontiguous update with attached pointers"); + return; + } + void *devaddr = (void *) (n->tgt->tgt_start + n->tgt_offset + + cur_node.host_start + - n->host_start); + size_t tmp_size = 0; + void *tmp = NULL; + if ((kind & typemask) == GOMP_MAP_TO_GRID) + omp_target_memcpy_rect_worker (devaddr, hostaddrs[i], + desc->elemsize, desc->ndims, + desc->length, desc->stride, + desc->index, desc->index, + desc->dim, desc->dim, devicep, + NULL, &tmp_size, &tmp); + else + omp_target_memcpy_rect_worker (hostaddrs[i], devaddr, + desc->elemsize, desc->ndims, + desc->length, desc->stride, + desc->index, desc->index, + desc->dim, desc->dim, NULL, + devicep, &tmp_size, &tmp); + } + i++; + } + else if (sizes[i]) + { + cur_node.host_start = (uintptr_t) hostaddrs[i]; + cur_node.host_end = cur_node.host_start + sizes[i]; + splay_tree_key n = splay_tree_lookup (&devicep->mem_map, &cur_node); + if (n) + { + if (n->host_start > cur_node.host_start + || n->host_end < cur_node.host_end) + { + gomp_mutex_unlock (&devicep->lock); + gomp_fatal ("Trying to update [%p..%p) object when " + "only [%p..%p) is mapped", + (void *) cur_node.host_start, + (void *) cur_node.host_end, + (void *) n->host_start, + (void *) n->host_end); + } - if (n->aux && n->aux->attach_count) - { - uintptr_t addr = cur_node.host_start; - while (addr < cur_node.host_end) - { - /* We have to be careful not to overwrite still attached - pointers during host<->device updates. */ - size_t i = (addr - cur_node.host_start) / sizeof (void *); - if (n->aux->attach_count[i] == 0) - { - void *devaddr = (void *) (n->tgt->tgt_start - + n->tgt_offset - + addr - n->host_start); - if (GOMP_MAP_COPY_TO_P (kind & typemask)) - gomp_copy_host2dev (devicep, NULL, - devaddr, (void *) addr, - sizeof (void *), false, NULL); - if (GOMP_MAP_COPY_FROM_P (kind & typemask)) - gomp_copy_dev2host (devicep, NULL, - (void *) addr, devaddr, - sizeof (void *)); - } - addr += sizeof (void *); - } - } - else - { - void *hostaddr = (void *) cur_node.host_start; - void *devaddr = (void *) (n->tgt->tgt_start + n->tgt_offset - + cur_node.host_start - - n->host_start); - size_t size = cur_node.host_end - cur_node.host_start; + if (n->aux && n->aux->attach_count) + { + uintptr_t addr = cur_node.host_start; + while (addr < cur_node.host_end) + { + /* We have to be careful not to overwrite still attached + pointers during host<->device updates. */ + size_t i = (addr - cur_node.host_start) / sizeof (void *); + if (n->aux->attach_count[i] == 0) + { + void *devaddr = (void *) (n->tgt->tgt_start + + n->tgt_offset + + addr - n->host_start); + if (GOMP_MAP_COPY_TO_P (kind & typemask)) + gomp_copy_host2dev (devicep, NULL, + devaddr, (void *) addr, + sizeof (void *), false, NULL); + if (GOMP_MAP_COPY_FROM_P (kind & typemask)) + gomp_copy_dev2host (devicep, NULL, + (void *) addr, devaddr, + sizeof (void *)); + } + addr += sizeof (void *); + } + } + else + { + void *hostaddr = (void *) cur_node.host_start; + void *devaddr = (void *) (n->tgt->tgt_start + n->tgt_offset + + cur_node.host_start + - n->host_start); + size_t size = cur_node.host_end - cur_node.host_start; - if (GOMP_MAP_COPY_TO_P (kind & typemask)) - gomp_copy_host2dev (devicep, NULL, devaddr, hostaddr, size, - false, NULL); - if (GOMP_MAP_COPY_FROM_P (kind & typemask)) - gomp_copy_dev2host (devicep, NULL, hostaddr, devaddr, size); - } - } - else - { - int kind = get_kind (short_mapkind, kinds, i); + if (GOMP_MAP_COPY_TO_P (kind & typemask)) + gomp_copy_host2dev (devicep, NULL, devaddr, hostaddr, size, + false, NULL); + if (GOMP_MAP_COPY_FROM_P (kind & typemask)) + gomp_copy_dev2host (devicep, NULL, hostaddr, devaddr, size); + } + } + else + { + int kind = get_kind (short_mapkind, kinds, i); - if (GOMP_MAP_PRESENT_P (kind)) - { - /* We already looked up the memory region above and it - was missing. */ - gomp_mutex_unlock (&devicep->lock); + if (GOMP_MAP_PRESENT_P (kind)) + { + /* We already looked up the memory region above and it + was missing. */ + gomp_mutex_unlock (&devicep->lock); #ifdef HAVE_INTTYPES_H - gomp_fatal ("present clause: not present on the device " - "(addr: %p, size: %"PRIu64" (0x%"PRIx64"), " - "dev: %d)", (void *) hostaddrs[i], - (uint64_t) sizes[i], (uint64_t) sizes[i], - devicep->target_id); + gomp_fatal ("present clause: not present on the device " + "(addr: %p, size: %"PRIu64" (0x%"PRIx64"), " + "dev: %d)", (void *) hostaddrs[i], + (uint64_t) sizes[i], (uint64_t) sizes[i], + devicep->target_id); #else - gomp_fatal ("present clause: not present on the device " - "(addr: %p, size: %lu (0x%lx), dev: %d)", - (void *) hostaddrs[i], (unsigned long) sizes[i], - (unsigned long) sizes[i], devicep->target_id); + gomp_fatal ("present clause: not present on the device " + "(addr: %p, size: %lu (0x%lx), dev: %d)", + (void *) hostaddrs[i], (unsigned long) sizes[i], + (unsigned long) sizes[i], devicep->target_id); #endif - } - } - } + } + } + } + } gomp_mutex_unlock (&devicep->lock); } @@ -4952,6 +5001,7 @@ omp_target_memcpy_async (void *dst, const void *src, size_t length, static int omp_target_memcpy_rect_worker (void *dst, const void *src, size_t element_size, int num_dims, const size_t *volume, + const size_t *strides, const size_t *dst_offsets, const size_t *src_offsets, const size_t *dst_dimensions, @@ -4965,7 +5015,7 @@ omp_target_memcpy_rect_worker (void *dst, const void *src, size_t element_size, size_t j, dst_off, src_off, length; int i, ret; - if (num_dims == 1) + if (num_dims == 1 && (!strides || strides[0] == 1)) { if (__builtin_mul_overflow (element_size, volume[0], &length) || __builtin_mul_overflow (element_size, dst_offsets[0], &dst_off) @@ -5019,6 +5069,38 @@ omp_target_memcpy_rect_worker (void *dst, const void *src, size_t element_size, } return ret ? 0 : EINVAL; } + else if (num_dims == 1 && strides) + { + size_t stride; + + assert ((src_devicep == NULL || dst_devicep == NULL) + && (src_devicep != NULL || dst_devicep != NULL)); + + if (__builtin_mul_overflow (element_size, dst_offsets[0], &dst_off) + || __builtin_mul_overflow (element_size, src_offsets[0], &src_off)) + return EINVAL; + + if (strides + && __builtin_mul_overflow (element_size, strides[0], &stride)) + return EINVAL; + + for (i = 0, ret = 1; i < volume[0] && ret; i++) + { + if (src_devicep == NULL) + ret = dst_devicep->host2dev_func (dst_devicep->target_id, + (char *) dst + dst_off, + (const char *) src + src_off, + element_size); + else if (dst_devicep == NULL) + ret = src_devicep->dev2host_func (src_devicep->target_id, + (char *) dst + dst_off, + (const char *) src + src_off, + element_size); + dst_off += stride; + src_off += stride; + } + return ret ? 0 : EINVAL; + } /* host->device, device->host and intra device. */ if (num_dims == 2 @@ -5083,13 +5165,19 @@ omp_target_memcpy_rect_worker (void *dst, const void *src, size_t element_size, if (__builtin_mul_overflow (dst_slice, dst_offsets[0], &dst_off) || __builtin_mul_overflow (src_slice, src_offsets[0], &src_off)) return EINVAL; + if (strides + && (__builtin_mul_overflow (dst_slice, strides[0], &dst_slice) + || __builtin_mul_overflow (src_slice, strides[0], &src_slice))) + return EINVAL; for (j = 0; j < volume[0]; j++) { ret = omp_target_memcpy_rect_worker ((char *) dst + dst_off, (const char *) src + src_off, element_size, num_dims - 1, - volume + 1, dst_offsets + 1, - src_offsets + 1, dst_dimensions + 1, + volume + 1, + strides ? strides + 1 : NULL, + dst_offsets + 1, src_offsets + 1, + dst_dimensions + 1, src_dimensions + 1, dst_devicep, src_devicep, tmp_size, tmp); if (ret) @@ -5139,7 +5227,7 @@ omp_target_memcpy_rect_copy (void *dst, const void *src, if (lock_dst) gomp_mutex_lock (&dst_devicep->lock); int ret = omp_target_memcpy_rect_worker (dst, src, element_size, num_dims, - volume, dst_offsets, src_offsets, + volume, NULL, dst_offsets, src_offsets, dst_dimensions, src_dimensions, dst_devicep, src_devicep, &tmp_size, &tmp); diff --git a/libgomp/testsuite/libgomp.c++/array-shaping-1.C b/libgomp/testsuite/libgomp.c++/array-shaping-1.C new file mode 100644 index 000000000000..6ff5f9475f6b --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/array-shaping-1.C @@ -0,0 +1,469 @@ +// { dg-do run { target offload_device_nonshared_as } } + +#include +#include + +volatile int yy = 4, zz = 2, str_str = 2; + +template +void foo() +{ + T *arr; + int x = 5; + T arr2d[10][10]; + + arr = new T[100]; + + /* Update whole reshaped array. */ + + memset (arr, 0, 100 * sizeof (T)); + +#pragma omp target enter data map(to: arr[:100]) + + for (int j = 0; j < x; j++) + for (int i = 0; i < 10; i++) + arr[j * 10 + i] = i ^ j; + +#pragma omp target update to(([10][x]) arr) + +#pragma omp target exit data map(from: arr[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + if (j < x) + assert (arr[j * 10 + i] == i ^ j); + else + assert (arr[j * 10 + i] == 0); + + + /* Strided update. */ + + memset (arr, 0, 100 * sizeof (T)); + +#pragma omp target enter data map(to: arr[:100]) + + for (int j = 0; j < 20; j++) + for (int i = 0; i < 5; i++) + arr[j * 5 + i] = i + j; + +#pragma omp target update to(([5][5]) arr[0:3][0:3:2]) + +#pragma omp target exit data map(from: arr[:100]) + + for (int j = 0; j < 20; j++) + for (int i = 0; i < 5; i++) + if (j < 3 && (i & 1) == 0 && i < 6) + assert (arr[j * 5 + i] == i + j); + else + assert (arr[j * 5 + i] == 0); + + + /* Reshaped update, contiguous. */ + + memset (arr, 0, 100 * sizeof (T)); + +#pragma omp target enter data map(to: arr[:100]) + + for (int j = 0; j < 20; j++) + for (int i = 0; i < 5; i++) + arr[j * 5 + i] = 2 * j + i; + +#pragma omp target update to(([5][5]) arr[0:5][0:5]) + +#pragma omp target exit data map(from: arr[:100]) + + for (int j = 0; j < 20; j++) + for (int i = 0; i < 5; i++) + if (j < 5 && i < 5) + assert (arr[j * 5 + i] == 2 * j + i); + else + assert (arr[j * 5 + i] == 0); + + + /* Strided update on actual array. */ + + memset (arr2d, 0, 100 * sizeof (T)); + +#pragma omp target enter data map(to: arr2d) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + arr2d[j][i] = j + 2 * i; + +#pragma omp target update to(arr2d[0:5:2][5:2]) + +#pragma omp target exit data map(from: arr2d) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + if ((j & 1) == 0 && i >= 5 && i < 7) + assert (arr2d[j][i] == j + 2 * i); + else + assert (arr2d[j][i] == 0); + + + /* Update with non-constant bounds. */ + + memset (arr, 0, 100 * sizeof (T)); + +#pragma omp target enter data map(to: arr[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + arr[j * 10 + i] = (2 * j) ^ i; + + x = 3; + int y = yy, z = zz, str = str_str; + /* This is actually [0:3:2] [4:2:2]. */ +#pragma omp target update to(([10][10]) arr[0:x:2][y:z:str]) + +#pragma omp target exit data map(from: arr[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + if ((j & 1) == 0 && j < 6 && (i & 1) == 0 && i >= 4 && i < 8) + assert (arr[j * 10 + i] == (2 * j) ^ i); + else + assert (arr[j * 10 + i] == 0); + + + /* Update with full "major" dimension. */ + + memset (arr, 0, 100 * sizeof (T)); + +#pragma omp target enter data map(to: arr[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + arr[j * 10 + i] = i + j; + +#pragma omp target update to(([10][10]) arr[0:10][3:1]) + +#pragma omp target exit data map(from: arr[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + if (i == 3) + assert (arr[j * 10 + i] == i + j); + else + assert (arr[j * 10 + i] == 0); + + + /* Update with full "minor" dimension. */ + + memset (arr, 0, 100 * sizeof (T)); + +#pragma omp target enter data map(to: arr[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + arr[j * 10 + i] = 3 * (i + j); + +#pragma omp target update to(([10][10]) arr[3:2][0:10]) + +#pragma omp target exit data map(from: arr[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + if (j >= 3 && j < 5) + assert (arr[j * 10 + i] == 3 * (i + j)); + else + assert (arr[j * 10 + i] == 0); + + + /* Rectangle update. */ + + memset (arr, 0, 100 * sizeof (T)); + +#pragma omp target enter data map(to: arr[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + arr[j * 10 + i] = 5 * (i + j); + +#pragma omp target update to(([10][10]) arr[3:2][0:9]) + +#pragma omp target exit data map(from: arr[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + if (j >= 3 && j < 5 && i < 9) + assert (arr[j * 10 + i] == 5 * (i + j)); + else + assert (arr[j * 10 + i] == 0); + + + /* One-dimensional strided update. */ + + memset (arr, 0, 100 * sizeof (T)); + +#pragma omp target enter data map(to: arr[:100]) + + for (int i = 0; i < 100; i++) + arr[i] = i + 99; + +#pragma omp target update to(([100]) arr[3:33:3]) + +#pragma omp target exit data map(from: arr[:100]) + + for (int i = 0; i < 100; i++) + if (i >= 3 && ((i - 3) % 3) == 0) + assert (arr[i] == i + 99); + else + assert (arr[i] == 0); + + + /* One-dimensional strided update without explicit array shape. */ + + memset (arr, 0, 100 * sizeof (T)); + +#pragma omp target enter data map(to: arr[:100]) + + for (int i = 0; i < 100; i++) + arr[i] = i + 121; + +#pragma omp target update to(arr[3:33:3]) + +#pragma omp target exit data map(from: arr[:100]) + + for (int i = 0; i < 100; i++) + if (i >= 3 && ((i - 3) % 3) == 0) + assert (arr[i] == i + 121); + else + assert (arr[i] == 0); + + delete[] arr; +} + +int main() +{ + int *arr; + int x = 5; + int arr2d[10][10]; + + arr = new int[100]; + + /* Update whole reshaped array. */ + + memset (arr, 0, 100 * sizeof (int)); + +#pragma omp target enter data map(to: arr[:100]) + + for (int j = 0; j < x; j++) + for (int i = 0; i < 10; i++) + arr[j * 10 + i] = i ^ j; + +#pragma omp target update to(([10][x]) arr) + +#pragma omp target exit data map(from: arr[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + if (j < x) + assert (arr[j * 10 + i] == i ^ j); + else + assert (arr[j * 10 + i] == 0); + + + /* Strided update. */ + + memset (arr, 0, 100 * sizeof (int)); + +#pragma omp target enter data map(to: arr[:100]) + + for (int j = 0; j < 20; j++) + for (int i = 0; i < 5; i++) + arr[j * 5 + i] = i + j; + +#pragma omp target update to(([5][5]) arr[0:3][0:3:2]) + +#pragma omp target exit data map(from: arr[:100]) + + for (int j = 0; j < 20; j++) + for (int i = 0; i < 5; i++) + if (j < 3 && (i & 1) == 0 && i < 6) + assert (arr[j * 5 + i] == i + j); + else + assert (arr[j * 5 + i] == 0); + + + /* Reshaped update, contiguous. */ + + memset (arr, 0, 100 * sizeof (int)); + +#pragma omp target enter data map(to: arr[:100]) + + for (int j = 0; j < 20; j++) + for (int i = 0; i < 5; i++) + arr[j * 5 + i] = 2 * j + i; + +#pragma omp target update to(([5][5]) arr[0:5][0:5]) + +#pragma omp target exit data map(from: arr[:100]) + + for (int j = 0; j < 20; j++) + for (int i = 0; i < 5; i++) + if (j < 5 && i < 5) + assert (arr[j * 5 + i] == 2 * j + i); + else + assert (arr[j * 5 + i] == 0); + + + /* Strided update on actual array. */ + + memset (arr2d, 0, 100 * sizeof (int)); + +#pragma omp target enter data map(to: arr2d) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + arr2d[j][i] = j + 2 * i; + +#pragma omp target update to(arr2d[0:5:2][5:2]) + +#pragma omp target exit data map(from: arr2d) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + if ((j & 1) == 0 && i >= 5 && i < 7) + assert (arr2d[j][i] == j + 2 * i); + else + assert (arr2d[j][i] == 0); + + + /* Update with non-constant bounds. */ + + memset (arr, 0, 100 * sizeof (int)); + +#pragma omp target enter data map(to: arr[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + arr[j * 10 + i] = (2 * j) ^ i; + + x = 3; + int y = yy, z = zz, str = str_str; + /* This is actually [0:3:2] [4:2:2]. */ +#pragma omp target update to(([10][10]) arr[0:x:2][y:z:str]) + +#pragma omp target exit data map(from: arr[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + if ((j & 1) == 0 && j < 6 && (i & 1) == 0 && i >= 4 && i < 8) + assert (arr[j * 10 + i] == (2 * j) ^ i); + else + assert (arr[j * 10 + i] == 0); + + + /* Update with full "major" dimension. */ + + memset (arr, 0, 100 * sizeof (int)); + +#pragma omp target enter data map(to: arr[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + arr[j * 10 + i] = i + j; + +#pragma omp target update to(([10][10]) arr[0:10][3:1]) + +#pragma omp target exit data map(from: arr[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + if (i == 3) + assert (arr[j * 10 + i] == i + j); + else + assert (arr[j * 10 + i] == 0); + + + /* Update with full "minor" dimension. */ + + memset (arr, 0, 100 * sizeof (int)); + +#pragma omp target enter data map(to: arr[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + arr[j * 10 + i] = 3 * (i + j); + +#pragma omp target update to(([10][10]) arr[3:2][0:10]) + +#pragma omp target exit data map(from: arr[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + if (j >= 3 && j < 5) + assert (arr[j * 10 + i] == 3 * (i + j)); + else + assert (arr[j * 10 + i] == 0); + + + /* Rectangle update. */ + + memset (arr, 0, 100 * sizeof (int)); + +#pragma omp target enter data map(to: arr[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + arr[j * 10 + i] = 5 * (i + j); + +#pragma omp target update to(([10][10]) arr[3:2][0:9]) + +#pragma omp target exit data map(from: arr[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + if (j >= 3 && j < 5 && i < 9) + assert (arr[j * 10 + i] == 5 * (i + j)); + else + assert (arr[j * 10 + i] == 0); + + + /* One-dimensional strided update. */ + + memset (arr, 0, 100 * sizeof (int)); + +#pragma omp target enter data map(to: arr[:100]) + + for (int i = 0; i < 100; i++) + arr[i] = i + 99; + +#pragma omp target update to(([100]) arr[3:33:3]) + +#pragma omp target exit data map(from: arr[:100]) + + for (int i = 0; i < 100; i++) + if (i >= 3 && ((i - 3) % 3) == 0) + assert (arr[i] == i + 99); + else + assert (arr[i] == 0); + + + /* One-dimensional strided update without explicit array shape. */ + + memset (arr, 0, 100 * sizeof (int)); + +#pragma omp target enter data map(to: arr[:100]) + + for (int i = 0; i < 100; i++) + arr[i] = i + 121; + +#pragma omp target update to(arr[3:33:3]) + +#pragma omp target exit data map(from: arr[:100]) + + for (int i = 0; i < 100; i++) + if (i >= 3 && ((i - 3) % 3) == 0) + assert (arr[i] == i + 121); + else + assert (arr[i] == 0); + + delete[] arr; + + foo (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/array-shaping-10.C b/libgomp/testsuite/libgomp.c++/array-shaping-10.C new file mode 100644 index 000000000000..648f02d34798 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/array-shaping-10.C @@ -0,0 +1,61 @@ +// { dg-do run { target offload_device_nonshared_as } } + +#include +#include + +#define N 10 + +template +void foo () +{ + T tarr[N * N]; + + memset (tarr, 0, N * N * sizeof (T)); + +#pragma omp target enter data map(to: tarr) + +#pragma omp target + { + for (int i = 0; i < N; i++) + for (int j = 0; j < N; j++) + tarr[i * N + j] = 2 * (i + j); + } + + /* An array, but cast to a pointer, then reshaped. */ +#pragma omp target update from(([N][N]) ((T *) &tarr[0])[4:3][5:3]) + + for (int i = 4; i < 7; i++) + for (int j = 5; j < 8; j++) + assert (tarr[i * N + j] == 2 * (i + j)); + +#pragma omp target exit data map(delete: tarr) +} + +int main () +{ + int iarr[N * N]; + + memset (iarr, 0, N * N * sizeof (int)); + +#pragma omp target enter data map(to: iarr) + +#pragma omp target + { + for (int i = 0; i < 10; i++) + for (int j = 0; j < 10; j++) + iarr[i * 10 + j] = i + j; + } + + /* An array, but cast to a pointer, then reshaped. */ +#pragma omp target update from(([10][10]) ((int *) &iarr[0])[4:3][4:3]) + + for (int i = 4; i < 7; i++) + for (int j = 4; j < 7; j++) + assert (iarr[i * 10 + j] == i + j); + +#pragma omp target exit data map(delete: iarr) + + foo (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/array-shaping-11.C b/libgomp/testsuite/libgomp.c++/array-shaping-11.C new file mode 100644 index 000000000000..6b15bd62fb1f --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/array-shaping-11.C @@ -0,0 +1,63 @@ +// { dg-do run { target offload_device_nonshared_as } } + +#include +#include + +#define N 10 + +template +void foo () +{ + T tarr_real[N * N]; + T (&tarr)[N * N] = tarr_real; + + memset (tarr, 0, N * N * sizeof (T)); + +#pragma omp target enter data map(to: tarr) + +#pragma omp target + { + for (int i = 0; i < N; i++) + for (int j = 0; j < N; j++) + tarr[i * N + j] = 2 * (i + j); + } + + /* A ref to an array, but cast to a pointer, then reshaped. */ +#pragma omp target update from(([N][N]) ((T *) &tarr[0])[4:3][5:3]) + + for (int i = 4; i < 7; i++) + for (int j = 5; j < 8; j++) + assert (tarr[i * N + j] == 2 * (i + j)); + +#pragma omp target exit data map(delete: tarr) +} + +int main () +{ + int iarr_real[N * N]; + int (&iarr)[N * N] = iarr_real; + + memset (iarr, 0, N * N * sizeof (int)); + +#pragma omp target enter data map(to: iarr) + +#pragma omp target + { + for (int i = 0; i < 10; i++) + for (int j = 0; j < 10; j++) + iarr[i * 10 + j] = i + j; + } + + /* A ref to an array, but cast to a pointer, then reshaped. */ +#pragma omp target update from(([10][10]) ((int *) &iarr[0])[4:3][4:3]) + + for (int i = 4; i < 7; i++) + for (int j = 4; j < 7; j++) + assert (iarr[i * 10 + j] == i + j); + +#pragma omp target exit data map(delete: iarr) + + foo (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/array-shaping-12.C b/libgomp/testsuite/libgomp.c++/array-shaping-12.C new file mode 100644 index 000000000000..103c99aa847e --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/array-shaping-12.C @@ -0,0 +1,65 @@ +// { dg-do run { target offload_device_nonshared_as } } + +#include +#include + +#define N 10 + +template +void foo () +{ + T tarr_real[N * N]; + T *tarrp = &tarr_real[0]; + T **tarrpp = &tarrp; + + memset (tarrp, 0, N * N * sizeof (T)); + +#pragma omp target enter data map(to: tarr_real) + +#pragma omp target + { + for (int i = 0; i < N; i++) + for (int j = 0; j < N; j++) + tarrp[i * N + j] = 2 * (i + j); + } + + /* A pointer with an extra indirection. */ +#pragma omp target update from(([N][N]) (*tarrpp)[4:3][5:3]) + + for (int i = 4; i < 7; i++) + for (int j = 5; j < 8; j++) + assert (tarrp[i * N + j] == 2 * (i + j)); + +#pragma omp target exit data map(delete: tarr_real) +} + +int main () +{ + int iarr_real[N * N]; + int *iarrp = &iarr_real[0]; + int **iarrpp = &iarrp; + + memset (iarrp, 0, N * N * sizeof (int)); + +#pragma omp target enter data map(to: iarr_real) + +#pragma omp target + { + for (int i = 0; i < 10; i++) + for (int j = 0; j < 10; j++) + iarrp[i * 10 + j] = i + j; + } + + /* A pointer with an extra indirection. */ +#pragma omp target update from(([10][10]) (*iarrpp)[4:3][4:3]) + + for (int i = 4; i < 7; i++) + for (int j = 4; j < 7; j++) + assert (iarrp[i * 10 + j] == i + j); + +#pragma omp target exit data map(delete: iarr_real) + + foo (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/array-shaping-13.C b/libgomp/testsuite/libgomp.c++/array-shaping-13.C new file mode 100644 index 000000000000..29345ca4264c --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/array-shaping-13.C @@ -0,0 +1,89 @@ +// { dg-do run { target offload_device_nonshared_as } } + +#include +#include + +#define N 10 + +template +void foo () +{ + T *tptr = new T[N * N * N]; + + memset (tptr, 0, N * N * N * sizeof (T)); + +#pragma omp target enter data map(to: tptr[0:N*N*N]) + +#pragma omp target + { + for (int i = 0; i < N; i++) + for (int j = 0; j < N; j++) + tptr[i * N * N + 4 * N + j] = 2 * (i + j); + } + + /* An array ref between two array sections. */ +#pragma omp target update from(([N][N][N]) tptr[4:3][4][5:3]) + + for (int i = 4; i < 7; i++) + for (int j = 5; j < 8; j++) + assert (tptr[i * N * N + 4 * N + j] == 2 * (i + j)); + + memset (tptr, 0, N * N * N * sizeof (T)); + + for (int i = 0; i < N; i++) + tptr[2 * N * N + i * N + 4] = 4 * i; + + /* Array section between two array refs. */ +#pragma omp target update to(([N][N][N]) tptr[2][3:6][4]) + +#pragma omp target exit data map(from: tptr[0:N*N*N]) + + for (int i = 3; i < 9; i++) + assert (tptr[2 * N * N + i * N + 4] == 4 * i); + +#pragma omp target exit data map(delete: tptr[0:N*N*N]) + + delete[] tptr; +} + +int main () +{ + int *iptr = new int[N * N * N]; + + memset (iptr, 0, N * N * N * sizeof (int)); + +#pragma omp target enter data map(to: iptr[0:N*N*N]) + +#pragma omp target + { + for (int i = 0; i < N; i++) + for (int j = 0; j < N; j++) + iptr[i * N * N + 4 * N + j] = i + j; + } + + /* An array ref between two array sections. */ +#pragma omp target update from(([N][N][N]) iptr[2:3][4][6:3]) + + for (int i = 2; i < 5; i++) + for (int j = 6; j < 9; j++) + assert (iptr[i * N * N + 4 * N + j] == i + j); + + memset (iptr, 0, N * N * N * sizeof (int)); + + for (int i = 0; i < N; i++) + iptr[2 * N * N + i * N + 4] = 3 * i; + + /* Array section between two array refs. */ +#pragma omp target update to(([N][N][N]) iptr[2][3:6][4]) + +#pragma omp target exit data map(from: iptr[0:N*N*N]) + + for (int i = 3; i < 9; i++) + assert (iptr[2 * N * N + i * N + 4] == 3 * i); + + delete[] iptr; + + foo (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/array-shaping-2.C b/libgomp/testsuite/libgomp.c++/array-shaping-2.C new file mode 100644 index 000000000000..027543e8d297 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/array-shaping-2.C @@ -0,0 +1,38 @@ +// { dg-do run { target offload_device_nonshared_as } } + +#include +#include + +template +void foo (T *w) +{ + memset (w, 0, sizeof (T) * 100); + +#pragma omp target enter data map(to: w[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + w[j * 10 + i] = i + j; + +#pragma omp target update to(([10][10]) w[3:2][1:8]) + +#pragma omp target exit data map(from: w[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + if (j >= 3 && j < 5 && i >= 1 && i < 9) + assert (w[j * 10 + i] == i + j); + else + assert (w[j * 10 + i] == 0); +} + +int main() +{ + int *arr = new int[100]; + + foo (arr); + + delete[] arr; + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/array-shaping-3.C b/libgomp/testsuite/libgomp.c++/array-shaping-3.C new file mode 100644 index 000000000000..09ff04bc1145 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/array-shaping-3.C @@ -0,0 +1,38 @@ +// { dg-do run { target offload_device_nonshared_as } } + +#include +#include + +template +void foo (double *w) +{ + memset (w, 0, sizeof (double) * 100); + +#pragma omp target enter data map(to: w[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + w[j * 10 + i] = i * 3 + j * 2; + +#pragma omp target update to(([C][D]) w[3:2][1:8]) + +#pragma omp target exit data map(from: w[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + if (j >= 3 && j < 5 && i >= 1 && i < 9) + assert (w[j * 10 + i] == i * 3 + j * 2); + else + assert (w[j * 10 + i] == 0.0f); +} + +int main() +{ + double *arr = new double[100]; + + foo<10, 10> (arr); + + delete[] arr; + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/array-shaping-4.C b/libgomp/testsuite/libgomp.c++/array-shaping-4.C new file mode 100644 index 000000000000..efa115e8be6b --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/array-shaping-4.C @@ -0,0 +1,38 @@ +// { dg-do run { target offload_device_nonshared_as } } + +#include +#include + +template +void foo (double *w) +{ + memset (w, 0, sizeof (double) * 100); + +#pragma omp target enter data map(to: w[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + w[j * 10 + i] = i * 2 + j * 3; + +#pragma omp target update to(([C][D]) w[3:2][1:8]) + +#pragma omp target exit data map(from: w[:100]) + + for (int j = 0; j < 10; j++) + for (int i = 0; i < 10; i++) + if (j >= 3 && j < 5 && i >= 1 && i < 9) + assert (w[j * 10 + i] == i * 2 + j * 3); + else + assert (w[j * 10 + i] == 0.0f); +} + +int main() +{ + double *arr = new double[100]; + + foo<10, 10> (arr); + + delete[] arr; + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/array-shaping-5.C b/libgomp/testsuite/libgomp.c++/array-shaping-5.C new file mode 100644 index 000000000000..7046a13c106f --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/array-shaping-5.C @@ -0,0 +1,38 @@ +// { dg-do run { target offload_device_nonshared_as } } + +#include +#include + +template +void foo (T *w, int e, int f, int g) +{ + memset (w, 0, sizeof (T) * 100); + +#pragma omp target enter data map(to: w[:100]) + + for (int j = 0; j < e; j++) + for (int i = 0; i < C; i++) + w[j * C + i] = i + j; + +#pragma omp target update to(([e][C]) w[3:2][f:g]) + +#pragma omp target exit data map(from: w[:100]) + + for (int j = 0; j < e; j++) + for (int i = 0; i < C; i++) + if (j >= 3 && j < 5 && i >= f && i < f + g) + assert (w[j * C + i] == i + j); + else + assert (w[j * C + i] == 0.0f); +} + +int main() +{ + float *arr = new float[100]; + + foo (arr, 10, 1, 8); + + delete[] arr; + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/array-shaping-6.C b/libgomp/testsuite/libgomp.c++/array-shaping-6.C new file mode 100644 index 000000000000..b960b5e58e14 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/array-shaping-6.C @@ -0,0 +1,54 @@ +// { dg-do run { target offload_device_nonshared_as } } + +#include +#include + +template +void foo (T *&aref) +{ +#pragma omp target enter data map(to: aref[:100]) + +#pragma omp target + { + for (int i = 0; i < 10; i++) + for (int j = 0; j < 10; j++) + aref[i * 10 + j] = i + j; + } + +#pragma omp target update from(([10][10]) aref[2:3:2][7:3]) + + for (int i = 2; i < 8; i += 2) + for (int j = 7; j < 10; j++) + assert (aref[i * 10 + j] == i + j); + +#pragma omp target exit data map(delete: aref[:100]) +} + +int main() +{ + float *arr = new float[100]; + float *&w = arr; + + memset (arr, 0, 100 * sizeof (float)); + +#pragma omp target enter data map(to: w[:100]) + +#pragma omp target + { + for (int i = 0; i < 10; i++) + for (int j = 0; j < 10; j++) + w[i * 10 + j] = i + j; + } + +#pragma omp target update from(([10][10]) w[4:3][4:3]) + + for (int i = 4; i < 7; i++) + for (int j = 4; j < 7; j++) + assert (w[i * 10 + j] == i + j); + +#pragma omp target exit data map(delete: w[:100]) + + foo (arr); + + delete[] arr; +} diff --git a/libgomp/testsuite/libgomp.c++/array-shaping-7.C b/libgomp/testsuite/libgomp.c++/array-shaping-7.C new file mode 100644 index 000000000000..b6193f8d619e --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/array-shaping-7.C @@ -0,0 +1,54 @@ +// { dg-do run { target offload_device_nonshared_as } } + +#include +#include + +template +void foo (T (&aref)[10][10]) +{ +#pragma omp target enter data map(to: aref) + +#pragma omp target + { + for (int i = 0; i < 10; i++) + for (int j = 0; j < 10; j++) + aref[i][j] = i + j; + } + +#pragma omp target update from(aref[2:3:2][7:3]) + + for (int i = 2; i < 8; i += 2) + for (int j = 7; j < 10; j++) + assert (aref[i][j] == i + j); + +#pragma omp target exit data map(delete: aref) +} + +int main() +{ + float arr2d[10][10]; + float (&w)[10][10] = arr2d; + + memset (&arr2d, 0, 100 * sizeof (float)); + +#pragma omp target enter data map(to: w) + +#pragma omp target + { + for (int i = 0; i < 10; i++) + for (int j = 0; j < 10; j++) + w[i][j] = i + j; + } + +#pragma omp target update from(w[4:3][4:3]) + + for (int i = 4; i < 7; i++) + for (int j = 4; j < 7; j++) + assert (w[i][j] == i + j); + +#pragma omp target exit data map(delete: w) + + foo (arr2d); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/array-shaping-8.C b/libgomp/testsuite/libgomp.c++/array-shaping-8.C new file mode 100644 index 000000000000..a96cf3cffb80 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/array-shaping-8.C @@ -0,0 +1,65 @@ +// { dg-do run { target offload_device_nonshared_as } } + +#include +#include + +template +struct C { + T *&aptr; + + C(T *&aptr_1) : aptr(aptr_1) + { + } +}; + +template +void foo (T *c) +{ +#pragma omp target enter data map(to: c->aptr, c->aptr[:100]) + +#pragma omp target + { + for (int i = 0; i < 10; i++) + for (int j = 0; j < 10; j++) + c->aptr[i * 10 + j] = i + j; + } + +#pragma omp target update from(([10][10]) c->aptr[2:3:2][7:3]) + + for (int i = 2; i < 8; i += 2) + for (int j = 7; j < 10; j++) + assert (c->aptr[i * 10 + j] == i + j); + +#pragma omp target exit data map(delete: c->aptr, c->aptr[:100]) +} + +int main() +{ + float *arr = new float[100]; + C cvar(arr); + + memset (arr, 0, 100 * sizeof (float)); + +#pragma omp target enter data map(to: cvar.aptr, cvar.aptr[:100]) + +#pragma omp target + { + for (int i = 0; i < 10; i++) + for (int j = 0; j < 10; j++) + cvar.aptr[i * 10 + j] = i + j; + } + +#pragma omp target update from(([10][10]) cvar.aptr[4:3][4:3]) + + for (int i = 4; i < 7; i++) + for (int j = 4; j < 7; j++) + assert (cvar.aptr[i * 10 + j] == i + j); + +#pragma omp target exit data map(delete: cvar.aptr, cvar.aptr[:100]) + + foo > (&cvar); + + delete[] arr; + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/array-shaping-9.C b/libgomp/testsuite/libgomp.c++/array-shaping-9.C new file mode 100644 index 000000000000..786fe9d11edb --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/array-shaping-9.C @@ -0,0 +1,95 @@ +// { dg-do run { target offload_device_nonshared_as } } + +#include +#include + +#define N 10 + +struct B { + int (&aref)[N][N]; + + B(int (&aref1)[N][N]) : aref(aref1) + { + } +}; + +template +struct C { + T (&aref)[S][S]; + + C(T (&aref1)[S][S]) : aref(aref1) + { + } +}; + +template +void foo (T *c) +{ +#pragma omp target enter data map(to: c->aref) + +#pragma omp target + { + for (int i = 0; i < 10; i++) + for (int j = 0; j < 10; j++) + c->aref[i][j] = 2 * (i + j); + } + +#pragma omp target update from(c->aref[2:3:2][7:3]) + + for (int i = 2; i < 8; i += 2) + for (int j = 7; j < 10; j++) + assert (c->aref[i][j] == 2 * (i + j)); + +#pragma omp target exit data map(delete: c->aref) +} + +int main() +{ + int iarr[N][N]; + float farr[N][N]; + B bvar(iarr); + C cvar(farr); + + memset (iarr, 0, N * N * sizeof (int)); + memset (farr, 0, N * N * sizeof (float)); + +#pragma omp target enter data map(to: bvar.aref) + +#pragma omp target + { + for (int i = 0; i < 10; i++) + for (int j = 0; j < 10; j++) + bvar.aref[i][j] = i + j; + } + +#pragma omp target update from(bvar.aref[4:3][4:3]) + + for (int i = 4; i < 7; i++) + for (int j = 4; j < 7; j++) + assert (bvar.aref[i][j] == i + j); + +#pragma omp target exit data map(delete: bvar.aref) + +#pragma omp target enter data map(to: cvar.aref) + +#pragma omp target + { + for (int i = 0; i < 10; i++) + for (int j = 0; j < 10; j++) + cvar.aref[i][j] = i + j; + } + +#pragma omp target update from(cvar.aref[4:3][4:3]) + + for (int i = 4; i < 7; i++) + for (int j = 4; j < 7; j++) + assert (cvar.aref[i][j] == i + j); + +#pragma omp target exit data map(delete: cvar.aref) + + memset (farr, 0, N * N * sizeof (float)); + + foo > (&cvar); + + return 0; +}