mirrors/gcc
0
0
Fork 0
mirror of https://gcc.gnu.org/git/gcc.git synced 2024-11-27 05:44:15 +08:00
Code Issues Actions Packages Projects Releases Wiki Activity

OpenMP: Support strided and shaped-array updates for C++

Browse Source 
This patch adds support for OpenMP 5.0 strided updates and the
array-shaping operator ("([x][y][z]) foo[0:n]...").  This is mostly for
C++ only so far, though necessary changes have been made to the C FE to
adjust for changes to shared data structures.

In terms of the implementation of various bits:

 - The OMP_ARRAY_SECTION tree code has been extended to take a 'stride'
   argument, and changes have been made throughout semantics.cc, etc. to
   take the new field into account -- including bounds checking.

 - A new type of cast operator has been added to represent the OpenMP
   array-shaping operator: OMP_ARRAYSHAPE_CAST_EXPR (1).

 - The address tokenization mechanism from previous patches has been
   extended with two new access kinds to represent noncontiguous array
   updates.

 - New mapping kinds have been added to represent noncontiguous updates:
   those which may be subject to array shaping, or have non-unit strides.
   These are processed by omp-low.cc into a kind of descriptor that is
   passed to the libgomp runtime (2).

The current patch reuses an extended version of the helper code for
omp_target_memcpy_rect, which may generate very many small host-device or
device-host copies.  (The "descriptor" has also been designed so reusing
that functionality is relatively straightforward.)  Optimising those
multiple copies, e.g. by packing them into a single transfer when it
would be beneficial, is left as the subject of a future patch.

This patch has some adjustments to the omp-low.cc code after Chung-Lin's
patch "OpenMP 5.0: Allow multiple clauses mapping same variable"
(325f085897), relative to the version last
posted for mainline.

Notes:

(1) In a bit more detail: the array-shaping operator has the same
precedence as a C-style cast, but applies to the whole expression,
including array-section specifiers. We parse it initially as if it
applies to the "value" of the whole expression:

  ([x][y]) ptr[0:10:2][1:5:2]

i.e., something like:

  ([x][y]) (ptr[0:10:2][1:5:2])

or as if the cast applies to the innermost/right-hand side array
section. Then, a little later in parsing (cp_parser_omp_var_list_no_open),
we rewrite it to apply to the inner pointer instead:

  (([x][y]) ptr)[0:10:2][1:5:2]

and that means a genuine multi-dimensional array or an array-shaped
pointer can be handled pretty much the same for the rest of
compilation. We use VIEW_CONVERT_EXPR for the "cast", unless we're
processing a template definition, where we use a new tree code instead.

(2) The new map kinds work like this. An update directive starts
out with OMP_CLAUSE_TO or OMP_CLAUSE_FROM clauses representing the
block in question and the direction of the needed transfer. If we
detect a noncontiguous update, we emit a list of mapping nodes (type
OMP_CLAUSE_MAP, with new kinds, so the "mapping group" machinery in
gimplify.cc can be reused):

  OMP_CLAUSE_TO -->

  GOMP_MAP_TO_GRID (VIEW_CONVERT_EXPR<int[x][y]>(ptr) [len: <element-size>])
  GOMP_MAP_GRID_DIM 0 [len: 10]   (i.e. [0:10:2])
  GOMP_MAP_GRID_STRIDE 2
  GOMP_MAP_GRID_DIM 1 [len: 5]    (i.e. [1:5:2])
  GOMP_MAP_GRID_STRIDE 2

During omp-low.cc, this sequence is reformulated into:

  GOMP_MAP_TO_GRID (ptr) [len: <whole array size>]
  GOMP_MAP_TO_PSET (&ptr_desc [len: <desc size>])

"ptr_desc" is a struct, stored statically or constructed on the (host)
stack, containing arrays representing the size of the whole array, the
rectangular subregion to transfer, and the stride with which to walk
over elements in each dimension.

2023-07-03  Julian Brown  <julian@codesourcery.com>

gcc/c-family/
	* c-common.h (expand_array_base): Update prototype.
	* c-omp.cc (c_omp_address_inspector::map_supported_p): Support
	VIEW_CONVERT_EXPR and ADDR_EXPR codes.
	(omp_expand_grid_dim): New function.
	(omp_handle_noncontig_array): New function.
	(c_omp_address_inspector:expand_array_base): Remove DECL_P parameter.
	Support noncontiguous array updates.
	(c_omp_address_inspector::expand_component_selector): Support
	noncontiguous array updates.
	(c_omp_address_inspector::expand_map_clause): Update calls to
	expand_array_base.
	* c-pretty-print.cc (c_pretty_printer::postfix_expression): Add
	OMP_ARRAY_SECTION stride support.

gcc/c/
	* c-parser.cc (c_parser_postfix_expression_after_primary): Dummy stride
	support (for now).
	(struct omp_dim): Add stride support.
	(c_parser_omp_variable_list): Likewise.
	* c-tree.h (build_omp_array_section): Update prototype.
	* c-typeck.cc (mark_exp_read): Add stride support for
	OMP_ARRAY_SECTION.
	(build_omp_array_section): Add stride support.
	(handle_omp_array_sections_1): Add minimal stride support.

gcc/cp/
	* cp-objcp-common.cc (cp_common_init_ts): Add array-shape cast
	support.
	* cp-tree.def (OMP_ARRAYSHAPE_CAST_EXPR): Add tree code.
	* cp-tree.h (DECLTYPE_FOR_OMP_ARRAYSHAPE_CAST): Add flag.
	(cp_omp_create_arrayshape_type, cp_build_omp_arrayshape_cast): Add
	prototypes.
	(grok_omp_array_section, build_omp_array_section): Add stride
	parameters.
	* decl.cc (create_anon_array_type): New function.
	(cp_omp_create_arrayshape_type): New function.
	* decl2.cc (grok_omp_array_section): Add stride parameter.
	(min_vis_expr_r): Add OMP_ARRAYSHAPE_CAST_EXPR support.
	* error.cc (dump_expr): Add stride support for OMP_ARRAY_SECTION.
	* mangle.cc (write_expression): Add OMP_ARRAYSHAPE_CAST_EXPR support.
	* operators.def (OMP_ARRAYSHAPE_CAST_EXPR): Add.
	* parser.cc (cp_parser_new): Initialise omp_array_shaping_op_p and
	omp_has_array_shape_p fields.
	(cp_parser_statement_expr): Don't allow array shaping op in statement
	exprs.
	(cp_parser_postfix_open_square_expression): Add stride parsing for
	array sections.  Use array section code to represent array refs if we
	have an array-shaping operator.
	(cp_parser_parenthesized_expression_list): Don't allow array-shaping
	op here.
	(cp_parser_cast_expression): Add array-shaping operator parsing.
	(cp_parser_lambda_expression): Don't allow array-shaping op in lambda
	body.
	(cp_parser_braced_list): Don't allow array-shaping op in braced list.
	(struct omp_dim): Add stride field.
	(cp_parser_var_list_no_open): Add stride/array shape support.
	(cp_parser_omp_target_update): Handle noncontiguous updates.
	* parser.h (cp_parser): Add omp_array_shaping_op_p and
	omp_has_array_shape_p fields.
	* pt.cc (tsubst): Add array-shape cast support.
	(tsubst_copy, tsubst_copy_and_build): Likewise. Add stride support for
	OMP_ARRAY_SECTION.
	(tsubst_omp_clause_decl): Add stride support for OMP_ARRAY_SECTION.
	* semantics.cc (handle_omp_array_sections_1): Add DISCONTIGUOUS
	parameter and stride support.
	(omp_array_section_low_bound): New function.
	(handle_omp_array_sections): Add DISCONTIGUOUS parameter and stride
	support.
	(finish_omp_clauses): Update calls to handle_omp_array_sections, and
	add noncontiguous array update support.
	(cp_build_omp_arrayshape_cast): New function.
	* typeck.cc (structural_comptypes): Add array-shape cast support.
	(build_omp_array_section): Add stride parameter.
	(check_for_casting_away_constness): Add OMP_ARRAYSHAPE_CAST_EXPR
	support.

gcc/
	* gimplify.cc (omp_group_last, omp_group_base): Add GOMP_MAP_TO_GRID,
	GOMP_MAP_FROM_GRID support.
	(gimplify_adjust_omp_clauses): Support new GOMP_MAP_GRID_DIM,
	GOMP_MAP_GRID_STRIDE mapping nodes.  Don't crash on e.g. misuse of
	ADDR_EXPR in mapping clauses.
	* omp-general.cc (omp_parse_noncontiguous_array): New function.
	(omp_parse_access_method): Add noncontiguous array support.
	(omp_parse_structure_base): Add array-shaping support.
	(debug_omp_tokenized_addr): Add ACCESS_NONCONTIG_ARRAY,
	ACCESS_NONCONTIG_REF_TO_ARRAY token support.
	* omp-general.h (access_method_kinds): Add ACCESS_NONCONTIG_ARRAY and
	ACCESS_NONCONTIG_REF_TO_ARRAY access kinds.
	* omp-low.cc (omp_noncontig_descriptor_type): New function.
	(scan_sharing_clauses): Support noncontiguous array updates.
	(lower_omp_target): Likewise.
	* tree-pretty-print.cc (dump_omp_clause): Add GOMP_MAP_TO_GRID,
	GOMP_MAP_FROM_GRID, GOMP_MAP_GRID_DIM, GOMP_MAP_GRID_STRIDE map kinds.
	(dump_generic_node): Add stride support for OMP_ARRAY_SECTION.
	* tree.def (OMP_ARRAY_SECTION): Add stride argument.

include/
	* gomp-constants.h (gomp_map_kind): Add GOMP_MAP_TO_GRID,
	GOMP_MAP_FROM_GRID, GOMP_MAP_GRID_DIM, GOMP_MAP_GRID_STRIDE map kinds.

gcc/testsuite/
	* g++.dg/gomp/array-shaping-1.C: New test.
	* g++.dg/gomp/array-shaping-2.C: New test.
	* g++.dg/gomp/bad-array-shaping-1.C: New test.
	* g++.dg/gomp/bad-array-shaping-2.C: New test.
	* g++.dg/gomp/bad-array-shaping-3.C: New test.
	* g++.dg/gomp/bad-array-shaping-4.C: New test.
	* g++.dg/gomp/bad-array-shaping-5.C: New test.
	* g++.dg/gomp/bad-array-shaping-6.C: New test.
	* g++.dg/gomp/bad-array-shaping-7.C: New test.
	* g++.dg/gomp/bad-array-shaping-8.C: New test.

libgomp/
	* libgomp.h (omp_noncontig_array_desc): New struct.
	* target.c (omp_target_memcpy_rect_worker): Add stride array
	parameter.  Forward declare.  Add STRIDES parameter and strided
	update support.
	(gomp_update): Add noncontiguous (strided/shaped) update support.
	* testsuite/libgomp.c++/array-shaping-1.C: New test.
	* testsuite/libgomp.c++/array-shaping-2.C: New test.
	* testsuite/libgomp.c++/array-shaping-3.C: New test.
	* testsuite/libgomp.c++/array-shaping-4.C: New test.
	* testsuite/libgomp.c++/array-shaping-5.C: New test.
	* testsuite/libgomp.c++/array-shaping-6.C: New test.
	* testsuite/libgomp.c++/array-shaping-7.C: New test.
	* testsuite/libgomp.c++/array-shaping-8.C: New test.
	* testsuite/libgomp.c++/array-shaping-9.C: New test.
	* testsuite/libgomp.c++/array-shaping-10.C: New test.
	* testsuite/libgomp.c++/array-shaping-11.C: New test.
	* testsuite/libgomp.c++/array-shaping-12.C: New test.
	* testsuite/libgomp.c++/array-shaping-13.C: New test.
This commit is contained in:
Julian Brown 2023-03-01 14:22:25 +00:00 committed by Paul-Antoine Arras
parent 61bc95e26d
commit 3d5bcb1a84
57 changed files with 3463 additions and 207 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

22
gcc/ChangeLog.omp
View File

@ -1,3 +1,25 @@
2023-07-03 Julian Brown <julian@codesourcery.com>
* gimplify.cc (omp_group_last, omp_group_base): Add GOMP_MAP_TO_GRID,
GOMP_MAP_FROM_GRID support.
(gimplify_adjust_omp_clauses): Support new GOMP_MAP_GRID_DIM,
GOMP_MAP_GRID_STRIDE mapping nodes. Don't crash on e.g. misuse of
ADDR_EXPR in mapping clauses.
* omp-general.cc (omp_parse_noncontiguous_array): New function.
(omp_parse_access_method): Add noncontiguous array support.
(omp_parse_structure_base): Add array-shaping support.
(debug_omp_tokenized_addr): Add ACCESS_NONCONTIG_ARRAY,
ACCESS_NONCONTIG_REF_TO_ARRAY token support.
* omp-general.h (access_method_kinds): Add ACCESS_NONCONTIG_ARRAY and
ACCESS_NONCONTIG_REF_TO_ARRAY access kinds.
* omp-low.cc (omp_noncontig_descriptor_type): New function.
(scan_sharing_clauses): Support noncontiguous array updates.
(lower_omp_target): Likewise.
* tree-pretty-print.cc (dump_omp_clause): Add GOMP_MAP_TO_GRID,
GOMP_MAP_FROM_GRID, GOMP_MAP_GRID_DIM, GOMP_MAP_GRID_STRIDE map kinds.
(dump_generic_node): Add stride support for OMP_ARRAY_SECTION.
* tree.def (OMP_ARRAY_SECTION): Add stride argument.
2023-06-30 Julian Brown <julian@codesourcery.com>
* gimplify.cc (dwarf2out.h): Include.

16
gcc/c-family/ChangeLog.omp
View File

@ -1,5 +1,21 @@
2023-09-05 Julian Brown <julian@codesourcery.com>
* c-common.h (expand_array_base): Update prototype.
* c-omp.cc (c_omp_address_inspector::map_supported_p): Support
VIEW_CONVERT_EXPR and ADDR_EXPR codes.
(omp_expand_grid_dim): New function.
(omp_handle_noncontig_array): New function.
(c_omp_address_inspector:expand_array_base): Remove DECL_P parameter.
Support noncontiguous array updates.
(c_omp_address_inspector::expand_component_selector): Support
noncontiguous array updates.
(c_omp_address_inspector::expand_map_clause): Update calls to
expand_array_base.
* c-pretty-print.cc (c_pretty_printer::postfix_expression): Add
OMP_ARRAY_SECTION stride support.
2023-07-03 Julian Brown <julian@codesourcery.com>
* c-common.h (expand_array_base, expand_component_selector,
expand_map_clause): Adjust member declarations.
* c-omp.cc (omp_expand_access_chain): Pass and return pointer to

198
gcc/c-family/c-omp.cc
View File

@ -3525,7 +3525,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
@ -3687,6 +3689,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 *
@ -3701,7 +3777,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)));
@ -3712,6 +3788,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);
@ -3725,7 +3802,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)
@ -3997,6 +4082,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;
@ -4048,8 +4167,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:
@ -4059,7 +4197,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));
@ -4068,7 +4206,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;
@ -4079,15 +4217,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,
@ -4104,12 +4242,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,
@ -4124,10 +4262,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);
@ -4142,13 +4280,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;
}
@ -4166,8 +4329,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);

5
gcc/c-family/c-pretty-print.cc
View File

@ -1658,6 +1658,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;

12
gcc/c/ChangeLog.omp
View File

@ -1,5 +1,17 @@
2023-09-05 Julian Brown <julian@codesourcery.com>
* c-parser.cc (c_parser_postfix_expression_after_primary): Dummy stride
support (for now).
(struct omp_dim): Add stride support.
(c_parser_omp_variable_list): Likewise.
* c-tree.h (build_omp_array_section): Update prototype.
* c-typeck.cc (mark_exp_read): Add stride support for
OMP_ARRAY_SECTION.
(build_omp_array_section): Add stride support.
(handle_omp_array_sections_1): Add minimal stride support.
2023-07-03 Julian Brown <julian@codesourcery.com>
* c-typeck.cc (handle_omp_array_sections): Pass pointer to clause to
process instead of clause.
(c_finish_omp_clauses): Update calls to handle_omp_array_sections.

32
gcc/c/c-parser.cc
View File

@ -12582,7 +12582,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);
@ -15211,11 +15211,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
@ -15345,7 +15345,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);
}
@ -15361,21 +15363,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)
{
@ -15384,7 +15387,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)
{
@ -15393,7 +15396,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))
@ -15548,7 +15552,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)
@ -15572,7 +15577,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)

2
gcc/c/c-tree.h
View File

@ -780,7 +780,7 @@ extern tree composite_type (tree, tree);
extern tree build_component_ref (location_t, tree, tree, location_t,
location_t);
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);

26
gcc/c/c-typeck.cc
View File

@ -2003,6 +2003,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;
@ -2927,7 +2929,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);
@ -2964,7 +2967,8 @@ build_omp_array_section (location_t loc, tree array, tree index, tree length)
sectype = 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);
}
@ -13863,7 +13867,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &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)
{
@ -13949,8 +13953,11 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &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)))
@ -13967,6 +13974,13 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &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))
@ -14183,7 +14197,9 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &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 (ort == C_ORT_ACC
&& OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)

52
gcc/cp/ChangeLog.omp
View File

@ -1,5 +1,57 @@
2023-09-05 Julian Brown <julian@codesourcery.com>
* cp-objcp-common.cc (cp_common_init_ts): Add array-shape cast
support.
* cp-tree.def (OMP_ARRAYSHAPE_CAST_EXPR): Add tree code.
* cp-tree.h (DECLTYPE_FOR_OMP_ARRAYSHAPE_CAST): Add flag.
(cp_omp_create_arrayshape_type, cp_build_omp_arrayshape_cast): Add
prototypes.
(grok_omp_array_section, build_omp_array_section): Add stride
parameters.
* decl.cc (create_anon_array_type): New function.
(cp_omp_create_arrayshape_type): New function.
* decl2.cc (grok_omp_array_section): Add stride parameter.
(min_vis_expr_r): Add OMP_ARRAYSHAPE_CAST_EXPR support.
* error.cc (dump_expr): Add stride support for OMP_ARRAY_SECTION.
* mangle.cc (write_expression): Add OMP_ARRAYSHAPE_CAST_EXPR support.
* operators.def (OMP_ARRAYSHAPE_CAST_EXPR): Add.
* parser.cc (cp_parser_new): Initialise omp_array_shaping_op_p and
omp_has_array_shape_p fields.
(cp_parser_statement_expr): Don't allow array shaping op in statement
exprs.
(cp_parser_postfix_open_square_expression): Add stride parsing for
array sections. Use array section code to represent array refs if we
have an array-shaping operator.
(cp_parser_parenthesized_expression_list): Don't allow array-shaping
op here.
(cp_parser_cast_expression): Add array-shaping operator parsing.
(cp_parser_lambda_expression): Don't allow array-shaping op in lambda
body.
(cp_parser_braced_list): Don't allow array-shaping op in braced list.
(struct omp_dim): Add stride field.
(cp_parser_var_list_no_open): Add stride/array shape support.
(cp_parser_omp_target_update): Handle noncontiguous updates.
* parser.h (cp_parser): Add omp_array_shaping_op_p and
omp_has_array_shape_p fields.
* pt.cc (tsubst): Add array-shape cast support.
(tsubst_copy, tsubst_copy_and_build): Likewise. Add stride support for
OMP_ARRAY_SECTION.
(tsubst_omp_clause_decl): Add stride support for OMP_ARRAY_SECTION.
* semantics.cc (handle_omp_array_sections_1): Add DISCONTIGUOUS
parameter and stride support.
(omp_array_section_low_bound): New function.
(handle_omp_array_sections): Add DISCONTIGUOUS parameter and stride
support.
(finish_omp_clauses): Update calls to handle_omp_array_sections, and
add noncontiguous array update support.
(cp_build_omp_arrayshape_cast): New function.
* typeck.cc (structural_comptypes): Add array-shape cast support.
(build_omp_array_section): Add stride parameter.
(check_for_casting_away_constness): Add OMP_ARRAYSHAPE_CAST_EXPR
support.
2023-07-03 Julian Brown <julian@codesourcery.com>
* semantics.cc (handle_omp_array_sections): Pass pointer to clause
instead of clause. Add PNEXT return parameter for next clause in list
to process.

1
gcc/cp/cp-objcp-common.cc
View File

@ -677,6 +677,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);

1
gcc/cp/cp-tree.def
View File

@ -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)

13
gcc/cp/cp-tree.h
View File

@ -511,6 +511,7 @@ extern GTY(()) tree cp_global_trees[CPTI_MAX];
OVL_LOOKUP_P (in OVERLOAD)
LOOKUP_FOUND_P (in RECORD_TYPE, UNION_TYPE, ENUMERAL_TYPE, NAMESPACE_DECL)
FNDECL_MANIFESTLY_CONST_EVALUATED (in FUNCTION_DECL)
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)
@ -4955,6 +4956,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
@ -7084,6 +7087,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<constructor_elt, va_gc> *);
extern tree cp_omp_create_arrayshape_type (location_t, tree,
vec<cp_expr> *);
extern tree check_var_type (tree, tree, location_t);
extern tree reshape_init (tree, tree, tsubst_flags_t);
extern tree next_aggregate_field (tree);
@ -7117,7 +7122,8 @@ extern void grokclassfn (tree, tree,
enum overload_flags);
extern tree grok_array_decl (location_t, tree, tree,
vec<tree, va_gc> **, 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);
@ -7980,6 +7986,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);
@ -8235,7 +8243,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);

75
gcc/cp/decl.cc
View File

@ -11875,6 +11875,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 %<auto%>");
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<cp_expr> *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

17
gcc/cp/decl2.cc
View File

@ -625,35 +625,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;
}
@ -2706,6 +2710,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:

5
gcc/cp/error.cc
View File

@ -2538,6 +2538,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;

1
gcc/cp/mangle.cc
View File

@ -3887,6 +3887,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;

1
gcc/cp/operators.def
View File

@ -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)

303
gcc/cp/parser.cc
View File

@ -4467,6 +4467,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;
@ -5471,6 +5477,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;
@ -8407,7 +8414,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)
@ -8440,9 +8447,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;
@ -8451,7 +8472,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;
@ -8459,11 +8480,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. */
@ -8785,6 +8818,7 @@ cp_parser_parenthesized_expression_list (cp_parser* parser,
vec<tree, va_gc> *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)
@ -8803,7 +8837,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);
@ -8872,6 +8908,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;
}
}
@ -8879,6 +8916,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;
}
@ -10147,6 +10185,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<cp_expr, 4> 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)
@ -10216,6 +10256,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
{
bool saved_in_type_id_in_expr_p = parser->in_type_id_in_expr_p;
@ -10237,6 +10299,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)
@ -10250,6 +10316,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)
@ -11390,6 +11474,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;
@ -11399,6 +11484,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;
/* The body of a lambda in a discarded statement is not discarded. */
bool discarded = in_discarded_stmt;
@ -11450,6 +11536,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 field is only used during parsing of the lambda. */
@ -26180,6 +26267,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;
@ -38152,11 +38240,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
@ -38189,10 +38277,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
@ -38203,49 +38303,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)
{
@ -38261,7 +38471,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)
{
@ -38270,7 +38480,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))
@ -38435,7 +38654,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
@ -38457,7 +38677,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;
@ -38470,6 +38691,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);
@ -47607,8 +47830,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 "

7
gcc/cp/parser.h
View File

@ -415,6 +415,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

32
gcc/cp/pt.cc
View File

@ -16959,6 +16959,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)
@ -17598,14 +17602,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. */
@ -20319,6 +20326,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),
@ -20504,7 +20519,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))
@ -20519,7 +20534,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:

262
gcc/cp/semantics.cc
View File

@ -5520,9 +5520,10 @@ public:
static tree
handle_omp_array_sections_1 (tree c, tree t, vec<tree> &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)
{
@ -5585,18 +5586,25 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &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)))
@ -5613,15 +5621,26 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &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))
@ -5632,9 +5651,15 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &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))
@ -5753,12 +5778,29 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &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)
@ -5766,7 +5808,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &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))
{
@ -5836,7 +5878,10 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &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 (ort == C_ORT_ACC
&& OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
@ -5857,10 +5902,15 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &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;
}
}
}
}
@ -5872,7 +5922,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &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)
@ -5891,15 +5941,45 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &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;
@ -5915,7 +5995,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)
@ -5957,6 +6037,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))
{
@ -5964,6 +6046,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
@ -5976,6 +6059,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;
@ -5985,10 +6073,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])
@ -6002,12 +6130,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
@ -6124,6 +6257,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;
@ -6135,9 +6271,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
@ -6164,6 +6297,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]". */
@ -6182,7 +6319,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
@ -7327,7 +7465,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;
@ -8463,7 +8601,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)
@ -8618,6 +8757,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:
@ -8632,8 +8774,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
{
@ -9226,7 +9371,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
{
@ -13664,6 +13809,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. */

12
gcc/cp/typeck.cc
View File

@ -1625,6 +1625,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;
@ -4804,7 +4807,7 @@ 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)
{
tree type = TREE_TYPE (array_expr);
gcc_assert (type);
@ -4843,8 +4846,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
@ -8213,6 +8216,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();
}

44
gcc/gimplify.cc
View File

@ -9745,6 +9745,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:
{
@ -9893,6 +9906,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);
@ -14864,7 +14881,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
@ -14954,6 +14973,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
@ -15046,8 +15079,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;
}

47
gcc/omp-general.cc
View File

@ -3575,6 +3575,32 @@ omp_parse_pointer (tree *expr0, bool *has_offset)
return false;
}
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_access_method (tree *expr0, enum access_method_kinds *kind)
{
@ -3583,6 +3609,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))
@ -3656,6 +3689,14 @@ omp_parse_structure_base (vec<omp_addr_token *> &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;
@ -3805,6 +3846,12 @@ debug_omp_tokenized_addr (vec<omp_addr_token *> &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:

4
gcc/omp-general.h
View File

@ -317,7 +317,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

402
gcc/omp-low.cc
View File

@ -1271,6 +1271,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. */
@ -1861,8 +1910,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)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
@ -2097,6 +2212,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
@ -13216,6 +13336,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:
@ -13243,6 +13367,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:
@ -13662,7 +13800,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<tree> 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++;
}
int tdim = tdims.length () - 1;
vec<constructor_elt, va_gc> *vdim;
vec<constructor_elt, va_gc> *vindex;
vec<constructor_elt, va_gc> *vlen;
vec<constructor_elt, va_gc> *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 >= 0)
{
/* 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<constructor_elt, va_gc> *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))

13
gcc/testsuite/ChangeLog.omp
View File

@ -1,3 +1,16 @@
2023-07-03 Julian Brown <julian@codesourcery.com>
* g++.dg/gomp/array-shaping-1.C: New test.
* g++.dg/gomp/array-shaping-2.C: New test.
* g++.dg/gomp/bad-array-shaping-1.C: New test.
* g++.dg/gomp/bad-array-shaping-2.C: New test.
* g++.dg/gomp/bad-array-shaping-3.C: New test.
* g++.dg/gomp/bad-array-shaping-4.C: New test.
* g++.dg/gomp/bad-array-shaping-5.C: New test.
* g++.dg/gomp/bad-array-shaping-6.C: New test.
* g++.dg/gomp/bad-array-shaping-7.C: New test.
* g++.dg/gomp/bad-array-shaping-8.C: New test.
2023-08-10 Julian Brown <julian@codesourcery.com>
* gfortran.dg/gomp/declare-mapper-31.f90: New test.

22
gcc/testsuite/g++.dg/gomp/array-shaping-1.C Normal file
View File

@ -0,0 +1,22 @@
// { dg-do compile }
// { dg-additional-options "-fdump-tree-original" }
template<typename T, typename E, int A, int B, int C, int D>
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<char, short, 3, 4, 5, 6> ();
return 0;
}

134
gcc/testsuite/g++.dg/gomp/array-shaping-2.C Normal file
View File

@ -0,0 +1,134 @@
// { dg-do compile }
// { dg-additional-options "-fdump-tree-original" }
template<typename T>
struct St
{
T ***ppptr;
T ***&rppptr;
St(T ***p, T ***&rp) : ppptr(p), rppptr(rp) { }
};
template<typename A, typename B>
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<int\[10\]>\(\*\*\*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<int\[10\]>\(\*\*\*\*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<B> *s = new St<B>(ppptr2, ppptr2);
St<B> **ps = &s;
St<B> **&rps = ps;
#pragma omp target update from(([10]) (**(*ps)->ppptr)[3:4:2])
// { dg-final { scan-tree-dump {map\(from_grid:VIEW_CONVERT_EXPR<long int\[10\]>\(\*\*\*\(\*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<long int\[10\]>\(\*\*\*\*\(\*\*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<char\[10\]>\(\*\*\*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<char\[10\]>\(\*\*\*\*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<short int\[10\]>\(\*\*\*\(\*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<short int\[10\]>\(\*\*\*\*\(\*\*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<int, long> ();
return 0;
}

47
gcc/testsuite/g++.dg/gomp/bad-array-shaping-1.C Normal file
View File

@ -0,0 +1,47 @@
// { dg-do compile }
#include <string.h>
#include <assert.h>
template<typename T, int C, int D>
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<float, 5, 20> (arr);
delete[] arr;
return 0;
}

52
gcc/testsuite/g++.dg/gomp/bad-array-shaping-2.C Normal file
View File

@ -0,0 +1,52 @@
// { dg-do compile }
#include <string.h>
#include <assert.h>
template<typename T, int C, int D>
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<float, 5, 20> (arr);
delete[] arr;
return 0;
}

53
gcc/testsuite/g++.dg/gomp/bad-array-shaping-3.C Normal file
View File

@ -0,0 +1,53 @@
// { dg-do compile }
#include <string.h>
#include <assert.h>
template<typename T>
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<float> (arr);
delete[] arr;
return 0;
}

60
gcc/testsuite/g++.dg/gomp/bad-array-shaping-4.C Normal file
View File

@ -0,0 +1,60 @@
// { dg-do compile }
#include <string.h>
#include <assert.h>
template<typename T>
extern T* baz(T*);
template<typename T>
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<float> (arr);
delete[] arr;
return 0;
}

55
gcc/testsuite/g++.dg/gomp/bad-array-shaping-5.C Normal file
View File

@ -0,0 +1,55 @@
// { dg-do compile }
// { dg-additional-options "-std=c++14" }
#include <string.h>
#include <assert.h>
template<typename T>
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<float> (arr);
delete[] arr;
return 0;
}

59
gcc/testsuite/g++.dg/gomp/bad-array-shaping-6.C Normal file
View File

@ -0,0 +1,59 @@
// { dg-do compile }
#include <string.h>
#include <assert.h>
template<typename T>
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<float> (arr);
delete[] arr;
return 0;
}

48
gcc/testsuite/g++.dg/gomp/bad-array-shaping-7.C Normal file
View File

@ -0,0 +1,48 @@
// { dg-do compile }
// { dg-additional-options "-std=c++11" }
#include <new>
template<typename T>
struct St {
T *pp;
};
template<typename T>
void foo (T *w)
{
alignas (St<T>) unsigned char buf[sizeof (St<T>)];
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<T> { ([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 {expected '\)' before 'sub1'} "" { target *-*-* } .-4 }
// { dg-error {expected an OpenMP clause before '\}' token} "" { target *-*-* } .-5 }
}
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 {expected '\)' before 'sub1'} "" { target *-*-* } .-4 }
// { dg-error {expected an OpenMP clause before '\}' token} "" { target *-*-* } .-5 }
// { dg-error {no match for 'operator\[\]'} "" { target *-*-* } .-6 }
// { dg-error {could not convert} "" { target *-*-* } .-7 }
// { dg-error {'#pragma omp target update' must contain at least one 'from' or 'to' clauses} "" { target *-*-* } .-8 }
return 0;
}

50
gcc/testsuite/g++.dg/gomp/bad-array-shaping-8.C Normal file
View File

@ -0,0 +1,50 @@
// { dg-do compile }
template<typename T>
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<char> ();
return 0;
}

17
gcc/tree-pretty-print.cc
View File

@ -1072,6 +1072,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;
@ -2852,6 +2864,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;

2
gcc/tree.def
View File

@ -1386,7 +1386,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)
/* TRANSACTION_EXPR tree code.
Operand 0: BODY: contains body of the transaction. */

5
include/ChangeLog.omp
View File

@ -1,3 +1,8 @@
2023-07-03 Julian Brown <julian@codesourcery.com>
* gomp-constants.h (gomp_map_kind): Add GOMP_MAP_TO_GRID,
GOMP_MAP_FROM_GRID, GOMP_MAP_GRID_DIM, GOMP_MAP_GRID_STRIDE map kinds.
2023-06-30 Julian Brown <julian@codesourcery.com>
* gomp-constants.h (gomp_map_kind): Add GOMP_MAP_MAPPING_GROUP.

7
include/gomp-constants.h
View File

@ -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) \

21
libgomp/ChangeLog.omp
View File

@ -1,3 +1,24 @@
2023-07-03 Julian Brown <julian@codesourcery.com>
* libgomp.h (omp_noncontig_array_desc): New struct.
* target.c (omp_target_memcpy_rect_worker): Add stride array
parameter. Forward declare. Add STRIDES parameter and strided
update support.
(gomp_update): Add noncontiguous (strided/shaped) update support.
* testsuite/libgomp.c++/array-shaping-1.C: New test.
* testsuite/libgomp.c++/array-shaping-2.C: New test.
* testsuite/libgomp.c++/array-shaping-3.C: New test.
* testsuite/libgomp.c++/array-shaping-4.C: New test.
* testsuite/libgomp.c++/array-shaping-5.C: New test.
* testsuite/libgomp.c++/array-shaping-6.C: New test.
* testsuite/libgomp.c++/array-shaping-7.C: New test.
* testsuite/libgomp.c++/array-shaping-8.C: New test.
* testsuite/libgomp.c++/array-shaping-9.C: New test.
* testsuite/libgomp.c++/array-shaping-10.C: New test.
* testsuite/libgomp.c++/array-shaping-11.C: New test.
* testsuite/libgomp.c++/array-shaping-12.C: New test.
* testsuite/libgomp.c++/array-shaping-13.C: New test.
2023-08-10 Julian Brown <julian@codesourcery.com>
* testsuite/libgomp.fortran/declare-mapper-30.f90: New test.

14
libgomp/libgomp.h
View File

@ -1326,6 +1326,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. */

252
libgomp/target.c
View File

@ -2311,6 +2311,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)
@ -2333,90 +2341,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);
}
@ -4917,6 +4966,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,
@ -4930,7 +4980,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)
@ -4984,6 +5034,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
@ -5048,13 +5130,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)
@ -5104,7 +5192,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);

469
libgomp/testsuite/libgomp.c++/array-shaping-1.C Normal file
View File

@ -0,0 +1,469 @@
// { dg-do run { target offload_device_nonshared_as } }
#include <string.h>
#include <assert.h>
volatile int yy = 4, zz = 2, str_str = 2;
template<typename T>
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<long> ();
return 0;
}

61
libgomp/testsuite/libgomp.c++/array-shaping-10.C Normal file
View File

@ -0,0 +1,61 @@
// { dg-do run { target offload_device_nonshared_as } }
#include <assert.h>
#include <string.h>
#define N 10
template<typename T>
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<unsigned short> ();
return 0;
}

63
libgomp/testsuite/libgomp.c++/array-shaping-11.C Normal file
View File

@ -0,0 +1,63 @@
// { dg-do run { target offload_device_nonshared_as } }
#include <assert.h>
#include <string.h>
#define N 10
template<typename T>
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<unsigned short> ();
return 0;
}

65
libgomp/testsuite/libgomp.c++/array-shaping-12.C Normal file
View File

@ -0,0 +1,65 @@
// { dg-do run { target offload_device_nonshared_as } }
#include <assert.h>
#include <string.h>
#define N 10
template<typename T>
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<unsigned short> ();
return 0;
}

89
libgomp/testsuite/libgomp.c++/array-shaping-13.C Normal file
View File

@ -0,0 +1,89 @@
// { dg-do run { target offload_device_nonshared_as } }
#include <assert.h>
#include <string.h>
#define N 10
template<typename T>
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<unsigned long> ();
return 0;
}

38
libgomp/testsuite/libgomp.c++/array-shaping-2.C Normal file
View File

@ -0,0 +1,38 @@
// { dg-do run { target offload_device_nonshared_as } }
#include <string.h>
#include <assert.h>
template<typename T>
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<int> (arr);
delete[] arr;
return 0;
}

38
libgomp/testsuite/libgomp.c++/array-shaping-3.C Normal file
View File

@ -0,0 +1,38 @@
// { dg-do run { target offload_device_nonshared_as } }
#include <string.h>
#include <assert.h>
template<int C, int D>
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;
}

38
libgomp/testsuite/libgomp.c++/array-shaping-4.C Normal file
View File

@ -0,0 +1,38 @@
// { dg-do run { target offload_device_nonshared_as } }
#include <string.h>
#include <assert.h>
template<auto C, auto D>
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;
}

38
libgomp/testsuite/libgomp.c++/array-shaping-5.C Normal file
View File

@ -0,0 +1,38 @@
// { dg-do run { target offload_device_nonshared_as } }
#include <string.h>
#include <assert.h>
template<typename T, auto C>
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<float, 10> (arr, 10, 1, 8);
delete[] arr;
return 0;
}

54
libgomp/testsuite/libgomp.c++/array-shaping-6.C Normal file
View File

@ -0,0 +1,54 @@
// { dg-do run { target offload_device_nonshared_as } }
#include <assert.h>
#include <string.h>
template<typename T>
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<float> (arr);
delete[] arr;
}

54
libgomp/testsuite/libgomp.c++/array-shaping-7.C Normal file
View File

@ -0,0 +1,54 @@
// { dg-do run { target offload_device_nonshared_as } }
#include <assert.h>
#include <string.h>
template<typename T>
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<float> (arr2d);
return 0;
}

65
libgomp/testsuite/libgomp.c++/array-shaping-8.C Normal file
View File

@ -0,0 +1,65 @@
// { dg-do run { target offload_device_nonshared_as } }
#include <assert.h>
#include <string.h>
template<typename T>
struct C {
T *&aptr;
C(T *&aptr_1) : aptr(aptr_1)
{
}
};
template<typename T>
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<float> 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<C<float> > (&cvar);
delete[] arr;
return 0;
}

95
libgomp/testsuite/libgomp.c++/array-shaping-9.C Normal file
View File

@ -0,0 +1,95 @@
// { dg-do run { target offload_device_nonshared_as } }
#include <assert.h>
#include <string.h>
#define N 10
struct B {
int (&aref)[N][N];
B(int (&aref1)[N][N]) : aref(aref1)
{
}
};
template<typename T, int S>
struct C {
T (&aref)[S][S];
C(T (&aref1)[S][S]) : aref(aref1)
{
}
};
template<typename T>
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<float, N> 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<C<float, N> > (&cvar);
return 0;
}