* [PATCH, OpenACC, v3] Non-contiguous array support for OpenACC data clauses
@ 2019-11-26 14:57 Chung-Lin Tang
2020-04-07 13:50 ` (v4 update) " Chung-Lin Tang
` (3 more replies)
0 siblings, 4 replies; 6+ messages in thread
From: Chung-Lin Tang @ 2019-11-26 14:57 UTC (permalink / raw)
To: gcc-patches, Thomas Schwinge; +Cc: Jakub Jelinek
[-- Attachment #1: Type: text/plain, Size: 5011 bytes --]
Hi Thomas,
this is a reorg of the last non-contiguous arrays patch. You'll notice that:
(1) A large part of the code has been pulled into oacc-parallel.c, with most
of the data structure declarations in oacc-int.h.
(2) target.c only contains relatively little code from gomp_map_vars_internal
that processes what GOACC_parallel_keyed/data_start gives it.
(3) Instead of directly passed in the map pointer, the array descriptor
pointers are now passed to GOACC_parallel_keyed/data_start using varargs.
(I believe the adding of '...' to GOACC_data_start does not break any
compatiblity)
(4) Along the way, I've added a 'gomp_map_vars_openacc' for specializing our
uses, which should shave off quite some code through inlining.
The GOMP_MAP_NONCONTIG_ARRAY_P maps are still placed at the beginning of the
recieved map sequence in this patch. It should still be relatively easy to
use a GOACC_FLAG_* to do so if deemed better before committing.
Thanks,
Chung-Lin
PR other/76739
gcc/c/
* c-typeck.c (handle_omp_array_sections_1): Add 'bool &non_contiguous'
parameter, adjust recursive call site, add cases for allowing
pointer based multi-dimensional arrays for OpenACC.
(handle_omp_array_sections): Adjust handle_omp_array_sections_1 call,
handle non-contiguous case to create dynamic array map.
gcc/cp/
* semantics.c (handle_omp_array_sections_1): Add 'bool &non_contiguous'
parameter, adjust recursive call site, add cases for allowing
pointer based multi-dimensional arrays for OpenACC.
(handle_omp_array_sections): Adjust handle_omp_array_sections_1 call,
handle non-contiguous case to create dynamic array map.
gcc/fortran/
* f95-lang.c (DEF_FUNCTION_TYPE_VAR_5): New symbol.
* types.def (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR): New type.
gcc/
* builtin-types.def (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR): New type.
* omp-builtins.def (BUILT_IN_GOACC_DATA_START): Adjust function type
to new BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR.
* gimplify.c (gimplify_scan_omp_clauses): Skip gimplification of
OMP_CLAUSE_SIZE of non-contiguous array maps (which is a TREE_LIST).
* omp-expand.c (expand_omp_target): Add non-contiguous array descriptor
pointers to variadic arguments.
* omp-low.c (append_field_to_record_type): New function.
(create_noncontig_array_descr_type): Likewise.
(create_noncontig_array_descr_init_code): Likewise.
(scan_sharing_clauses): For non-contiguous array map kinds, check for
supported dimension structure, and install non-contiguous array
variable into current omp_context.
(reorder_noncontig_array_clauses): New function.
(scan_omp_target): Call reorder_noncontig_array_clauses to place
non-contiguous array map clauses at beginning of clause sequence.
(lower_omp_target): Add handling for non-contiguous array map kinds,
add all created non-contiguous array descriptors to
gimple_omp_target_data_arg.
gcc/testsuite/
* c-c++-common/goacc/noncontig_array-1.c: New test.
libgomp/
* libgomp_g.h (GOACC_data_start): Add variadic '...' to declaration.
* libgomp.h (gomp_map_vars_openacc): New function declaration.
* oacc-int.h (struct goacc_ncarray_dim): New struct declaration.
(struct goacc_ncarray_descr_type): Likewise.
(struct goacc_ncarray): Likewise.
(struct goacc_ncarray_info): Likewise.
(goacc_noncontig_array_create_ptrblock): New function declaration.
* oacc-parallel.c (goacc_noncontig_array_count_rows): New function.
(goacc_noncontig_array_compute_sizes): Likewise.
(goacc_noncontig_array_fill_rows_1): Likewise.
(goacc_noncontig_array_fill_rows): Likewise.
(goacc_process_noncontiguous_arrays): Likewise.
(goacc_noncontig_array_create_ptrblock): Likewise.
(GOACC_parallel_keyed): Use goacc_process_noncontiguous_arrays to
handle non-contiguous array descriptors at end of varargs, adjust
to use gomp_map_vars_openacc.
(GOACC_data_start): Likewise. Adjust function type to accept varargs.
* target.c (gomp_map_vars_internal): Add struct goacc_ncarray_info *
nca_info parameter, add handling code for non-contiguous arrays.
(gomp_map_vars_openacc): Add new function for specialization of
gomp_map_vars_internal for OpenACC structured region usage.
* testsuite/libgomp.oacc-c-c++-common/noncontig_array-1.c: New test.
* testsuite/libgomp.oacc-c-c++-common/noncontig_array-2.c: New test.
* testsuite/libgomp.oacc-c-c++-common/noncontig_array-3.c: New test.
* testsuite/libgomp.oacc-c-c++-common/noncontig_array-4.c: New test.
* testsuite/libgomp.oacc-c-c++-common/noncontig_array-utils.h: Support
header for new tests.
[-- Attachment #2: openacc-noncontig-arrays-v3.patch --]
[-- Type: text/plain, Size: 59918 bytes --]
Index: include/gomp-constants.h
===================================================================
--- include/gomp-constants.h (revision 278656)
+++ include/gomp-constants.h (working copy)
@@ -40,6 +40,7 @@
#define GOMP_MAP_FLAG_SPECIAL_0 (1 << 2)
#define GOMP_MAP_FLAG_SPECIAL_1 (1 << 3)
#define GOMP_MAP_FLAG_SPECIAL_2 (1 << 4)
+#define GOMP_MAP_FLAG_SPECIAL_3 (1 << 5)
#define GOMP_MAP_FLAG_SPECIAL (GOMP_MAP_FLAG_SPECIAL_1 \
| GOMP_MAP_FLAG_SPECIAL_0)
/* Flag to force a specific behavior (or else, trigger a run-time error). */
@@ -127,6 +128,26 @@ enum gomp_map_kind
/* Decrement usage count and deallocate if zero. */
GOMP_MAP_RELEASE = (GOMP_MAP_FLAG_SPECIAL_2
| GOMP_MAP_DELETE),
+ /* Mapping kinds for non-contiguous arrays. */
+ GOMP_MAP_NONCONTIG_ARRAY = (GOMP_MAP_FLAG_SPECIAL_3),
+ GOMP_MAP_NONCONTIG_ARRAY_TO = (GOMP_MAP_NONCONTIG_ARRAY
+ | GOMP_MAP_TO),
+ GOMP_MAP_NONCONTIG_ARRAY_FROM = (GOMP_MAP_NONCONTIG_ARRAY
+ | GOMP_MAP_FROM),
+ GOMP_MAP_NONCONTIG_ARRAY_TOFROM = (GOMP_MAP_NONCONTIG_ARRAY
+ | GOMP_MAP_TOFROM),
+ GOMP_MAP_NONCONTIG_ARRAY_FORCE_TO = (GOMP_MAP_NONCONTIG_ARRAY_TO
+ | GOMP_MAP_FLAG_FORCE),
+ GOMP_MAP_NONCONTIG_ARRAY_FORCE_FROM = (GOMP_MAP_NONCONTIG_ARRAY_FROM
+ | GOMP_MAP_FLAG_FORCE),
+ GOMP_MAP_NONCONTIG_ARRAY_FORCE_TOFROM = (GOMP_MAP_NONCONTIG_ARRAY_TOFROM
+ | GOMP_MAP_FLAG_FORCE),
+ GOMP_MAP_NONCONTIG_ARRAY_ALLOC = (GOMP_MAP_NONCONTIG_ARRAY
+ | GOMP_MAP_ALLOC),
+ GOMP_MAP_NONCONTIG_ARRAY_FORCE_ALLOC = (GOMP_MAP_NONCONTIG_ARRAY
+ | GOMP_MAP_FORCE_ALLOC),
+ GOMP_MAP_NONCONTIG_ARRAY_FORCE_PRESENT = (GOMP_MAP_NONCONTIG_ARRAY
+ | GOMP_MAP_FORCE_PRESENT),
/* Internal to GCC, not used in libgomp. */
/* Do not map, but pointer assign a pointer instead. */
@@ -155,6 +176,8 @@ enum gomp_map_kind
#define GOMP_MAP_ALWAYS_P(X) \
(GOMP_MAP_ALWAYS_TO_P (X) || ((X) == GOMP_MAP_ALWAYS_FROM))
+#define GOMP_MAP_NONCONTIG_ARRAY_P(X) \
+ ((X) & GOMP_MAP_NONCONTIG_ARRAY)
/* Asynchronous behavior. Keep in sync with
libgomp/{openacc.h,openacc.f90,openacc_lib.h}:acc_async_t. */
Index: gcc/builtin-types.def
===================================================================
--- gcc/builtin-types.def (revision 278656)
+++ gcc/builtin-types.def (working copy)
@@ -822,6 +822,9 @@ DEF_FUNCTION_TYPE_VAR_5 (BT_FN_INT_STRING_SIZE_INT
DEF_FUNCTION_TYPE_VAR_5 (BT_FN_INT_INT_INT_INT_INT_INT_VAR,
BT_INT, BT_INT, BT_INT, BT_INT, BT_INT, BT_INT)
+DEF_FUNCTION_TYPE_VAR_5 (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR,
+ BT_VOID, BT_INT, BT_SIZE, BT_PTR, BT_PTR, BT_PTR)
+
DEF_FUNCTION_TYPE_VAR_6 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_VAR,
BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE,
BT_PTR, BT_PTR, BT_PTR)
Index: gcc/c/c-typeck.c
===================================================================
--- gcc/c/c-typeck.c (revision 278656)
+++ gcc/c/c-typeck.c (working copy)
@@ -12876,12 +12876,14 @@ c_finish_omp_cancellation_point (location_t loc, t
<= FIRST_NON_ONE we diagnose non-contiguous arrays if low bound isn't
0 or length isn't the array domain max + 1, for > FIRST_NON_ONE we
can if MAYBE_ZERO_LEN is false. MAYBE_ZERO_LEN will be true in the above
- case though, as some lengths could be zero. */
+ case though, as some lengths could be zero.
+ NON_CONTIGUOUS will be true if this is an OpenACC non-contiguous array
+ section. */
static tree
handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
bool &maybe_zero_len, unsigned int &first_non_one,
- enum c_omp_region_type ort)
+ bool &non_contiguous, enum c_omp_region_type ort)
{
tree ret, low_bound, length, type;
if (TREE_CODE (t) != TREE_LIST)
@@ -12966,7 +12968,8 @@ handle_omp_array_sections_1 (tree c, tree t, vec<t
}
ret = handle_omp_array_sections_1 (c, TREE_CHAIN (t), types,
- maybe_zero_len, first_non_one, ort);
+ maybe_zero_len, first_non_one,
+ non_contiguous, ort);
if (ret == error_mark_node || ret == NULL_TREE)
return ret;
@@ -13173,14 +13176,21 @@ handle_omp_array_sections_1 (tree c, tree t, vec<t
return error_mark_node;
}
/* If there is a pointer type anywhere but in the very first
- array-section-subscript, the array section can't be contiguous. */
+ array-section-subscript, the array section can't be contiguous.
+ Note that OpenACC does accept these kinds of non-contiguous pointer
+ based arrays. */
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND
&& TREE_CODE (TREE_CHAIN (t)) == TREE_LIST)
{
- 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 (ort == C_ORT_ACC)
+ non_contiguous = true;
+ 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;
+ }
}
}
else
@@ -13209,6 +13219,7 @@ handle_omp_array_sections (tree c, enum c_omp_regi
{
bool maybe_zero_len = false;
unsigned int first_non_one = 0;
+ bool non_contiguous = false;
auto_vec<tree, 10> types;
tree *tp = &OMP_CLAUSE_DECL (c);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND
@@ -13218,7 +13229,7 @@ handle_omp_array_sections (tree c, enum c_omp_regi
tp = &TREE_VALUE (*tp);
tree first = handle_omp_array_sections_1 (c, *tp, types,
maybe_zero_len, first_non_one,
- ort);
+ non_contiguous, ort);
if (first == error_mark_node)
return true;
if (first == NULL_TREE)
@@ -13251,6 +13262,7 @@ handle_omp_array_sections (tree c, enum c_omp_regi
unsigned int num = types.length (), i;
tree t, side_effects = NULL_TREE, size = NULL_TREE;
tree condition = NULL_TREE;
+ tree ncarray_dims = NULL_TREE;
if (int_size_in_bytes (TREE_TYPE (first)) <= 0)
maybe_zero_len = true;
@@ -13274,6 +13286,13 @@ handle_omp_array_sections (tree c, enum c_omp_regi
length = fold_convert (sizetype, length);
if (low_bound == NULL_TREE)
low_bound = integer_zero_node;
+
+ if (non_contiguous)
+ {
+ ncarray_dims = tree_cons (low_bound, length, ncarray_dims);
+ continue;
+ }
+
if (!maybe_zero_len && i > first_non_one)
{
if (integer_nonzerop (low_bound))
@@ -13370,6 +13389,14 @@ handle_omp_array_sections (tree c, enum c_omp_regi
size = size_binop (MULT_EXPR, size, l);
}
}
+ if (non_contiguous)
+ {
+ int kind = OMP_CLAUSE_MAP_KIND (c);
+ OMP_CLAUSE_SET_MAP_KIND (c, kind | GOMP_MAP_NONCONTIG_ARRAY);
+ OMP_CLAUSE_DECL (c) = t;
+ OMP_CLAUSE_SIZE (c) = ncarray_dims;
+ return false;
+ }
if (side_effects)
size = build2 (COMPOUND_EXPR, sizetype, side_effects, size);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION
Index: gcc/cp/semantics.c
===================================================================
--- gcc/cp/semantics.c (revision 278656)
+++ gcc/cp/semantics.c (working copy)
@@ -4735,12 +4735,14 @@ omp_privatize_field (tree t, bool shared)
<= FIRST_NON_ONE we diagnose non-contiguous arrays if low bound isn't
0 or length isn't the array domain max + 1, for > FIRST_NON_ONE we
can if MAYBE_ZERO_LEN is false. MAYBE_ZERO_LEN will be true in the above
- case though, as some lengths could be zero. */
+ case though, as some lengths could be zero.
+ NON_CONTIGUOUS will be true if this is an OpenACC non-contiguous array
+ section. */
static tree
handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
bool &maybe_zero_len, unsigned int &first_non_one,
- enum c_omp_region_type ort)
+ bool &non_contiguous, enum c_omp_region_type ort)
{
tree ret, low_bound, length, type;
if (TREE_CODE (t) != TREE_LIST)
@@ -4825,7 +4827,8 @@ handle_omp_array_sections_1 (tree c, tree t, vec<t
&& TREE_CODE (TREE_CHAIN (t)) == FIELD_DECL)
TREE_CHAIN (t) = omp_privatize_field (TREE_CHAIN (t), false);
ret = handle_omp_array_sections_1 (c, TREE_CHAIN (t), types,
- maybe_zero_len, first_non_one, ort);
+ maybe_zero_len, first_non_one,
+ non_contiguous, ort);
if (ret == error_mark_node || ret == NULL_TREE)
return ret;
@@ -5044,14 +5047,21 @@ handle_omp_array_sections_1 (tree c, tree t, vec<t
return error_mark_node;
}
/* If there is a pointer type anywhere but in the very first
- array-section-subscript, the array section can't be contiguous. */
+ array-section-subscript, the array section can't be contiguous.
+ Note that OpenACC does accept these kinds of non-contiguous pointer
+ based arrays. */
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND
&& TREE_CODE (TREE_CHAIN (t)) == TREE_LIST)
{
- 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 (ort == C_ORT_ACC)
+ non_contiguous = true;
+ 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;
+ }
}
}
else
@@ -5091,6 +5101,7 @@ handle_omp_array_sections (tree c, enum c_omp_regi
{
bool maybe_zero_len = false;
unsigned int first_non_one = 0;
+ bool non_contiguous = false;
auto_vec<tree, 10> types;
tree *tp = &OMP_CLAUSE_DECL (c);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND
@@ -5100,7 +5111,7 @@ handle_omp_array_sections (tree c, enum c_omp_regi
tp = &TREE_VALUE (*tp);
tree first = handle_omp_array_sections_1 (c, *tp, types,
maybe_zero_len, first_non_one,
- ort);
+ non_contiguous, ort);
if (first == error_mark_node)
return true;
if (first == NULL_TREE)
@@ -5134,6 +5145,7 @@ handle_omp_array_sections (tree c, enum c_omp_regi
unsigned int num = types.length (), i;
tree t, side_effects = NULL_TREE, size = NULL_TREE;
tree condition = NULL_TREE;
+ tree ncarray_dims = NULL_TREE;
if (int_size_in_bytes (TREE_TYPE (first)) <= 0)
maybe_zero_len = true;
@@ -5159,6 +5171,13 @@ handle_omp_array_sections (tree c, enum c_omp_regi
length = fold_convert (sizetype, length);
if (low_bound == NULL_TREE)
low_bound = integer_zero_node;
+
+ if (non_contiguous)
+ {
+ ncarray_dims = tree_cons (low_bound, length, ncarray_dims);
+ continue;
+ }
+
if (!maybe_zero_len && i > first_non_one)
{
if (integer_nonzerop (low_bound))
@@ -5250,6 +5269,14 @@ handle_omp_array_sections (tree c, enum c_omp_regi
}
if (!processing_template_decl)
{
+ if (non_contiguous)
+ {
+ int kind = OMP_CLAUSE_MAP_KIND (c);
+ OMP_CLAUSE_SET_MAP_KIND (c, kind | GOMP_MAP_NONCONTIG_ARRAY);
+ OMP_CLAUSE_DECL (c) = t;
+ OMP_CLAUSE_SIZE (c) = ncarray_dims;
+ return false;
+ }
if (side_effects)
size = build2 (COMPOUND_EXPR, sizetype, side_effects, size);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION
Index: gcc/fortran/f95-lang.c
===================================================================
--- gcc/fortran/f95-lang.c (revision 278656)
+++ gcc/fortran/f95-lang.c (working copy)
@@ -632,6 +632,8 @@ gfc_init_builtin_functions (void)
#define DEF_FUNCTION_TYPE_VAR_0(NAME, RETURN) NAME,
#define DEF_FUNCTION_TYPE_VAR_1(NAME, RETURN, ARG1) NAME,
#define DEF_FUNCTION_TYPE_VAR_2(NAME, RETURN, ARG1, ARG2) NAME,
+#define DEF_FUNCTION_TYPE_VAR_5(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5) \
+ NAME,
#define DEF_FUNCTION_TYPE_VAR_6(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
ARG6) NAME,
#define DEF_FUNCTION_TYPE_VAR_7(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
Index: gcc/fortran/types.def
===================================================================
--- gcc/fortran/types.def (revision 278656)
+++ gcc/fortran/types.def (working copy)
@@ -269,6 +269,9 @@ DEF_FUNCTION_TYPE_VAR_7 (BT_FN_VOID_INT_SIZE_PTR_P
BT_VOID, BT_INT, BT_SIZE, BT_PTR, BT_PTR,
BT_PTR, BT_INT, BT_INT)
+DEF_FUNCTION_TYPE_VAR_5 (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR,
+ BT_VOID, BT_INT, BT_SIZE, BT_PTR, BT_PTR, BT_PTR)
+
DEF_FUNCTION_TYPE_VAR_6 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_VAR,
BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE,
BT_PTR, BT_PTR, BT_PTR)
Index: gcc/gimplify.c
===================================================================
--- gcc/gimplify.c (revision 278656)
+++ gcc/gimplify.c (working copy)
@@ -8641,9 +8641,17 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_se
if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
OMP_CLAUSE_SIZE (c) = DECL_P (decl) ? DECL_SIZE_UNIT (decl)
: TYPE_SIZE_UNIT (TREE_TYPE (decl));
- if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p,
- NULL, is_gimple_val, fb_rvalue) == GS_ERROR)
+ if (GOMP_MAP_NONCONTIG_ARRAY_P (OMP_CLAUSE_MAP_KIND (c)))
{
+ gcc_assert (OMP_CLAUSE_SIZE (c)
+ && TREE_CODE (OMP_CLAUSE_SIZE (c)) == TREE_LIST);
+ /* For non-contiguous array maps, OMP_CLAUSE_SIZE is a TREE_LIST
+ of the individual array dimensions, which gimplify_expr doesn't
+ handle, so skip the call to gimplify_expr here. */
+ }
+ else if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p,
+ NULL, is_gimple_val, fb_rvalue) == GS_ERROR)
+ {
remove = true;
break;
}
Index: gcc/omp-builtins.def
===================================================================
--- gcc/omp-builtins.def (revision 278656)
+++ gcc/omp-builtins.def (working copy)
@@ -32,7 +32,7 @@ along with GCC; see the file COPYING3. If not see
DEF_GOACC_BUILTIN (BUILT_IN_ACC_GET_DEVICE_TYPE, "acc_get_device_type",
BT_FN_INT, ATTR_NOTHROW_LIST)
DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DATA_START, "GOACC_data_start",
- BT_FN_VOID_INT_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
+ BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR, ATTR_NOTHROW_LIST)
DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DATA_END, "GOACC_data_end",
BT_FN_VOID, ATTR_NOTHROW_LIST)
DEF_GOACC_BUILTIN (BUILT_IN_GOACC_ENTER_EXIT_DATA, "GOACC_enter_exit_data",
Index: gcc/omp-expand.c
===================================================================
--- gcc/omp-expand.c (revision 278656)
+++ gcc/omp-expand.c (working copy)
@@ -8458,6 +8458,21 @@ expand_omp_target (struct omp_region *region)
if (tagging)
/* Push terminal marker - zero. */
args.safe_push (oacc_launch_pack (0, NULL_TREE, 0));
+
+ /* We assume index >= 3 in gimple_omp_target_data_arg are non-contiguous
+ array descriptor pointer arguments. */
+ if (TREE_VEC_LENGTH (t) > 3
+ && (start_ix == BUILT_IN_GOACC_DATA_START
+ || start_ix == BUILT_IN_GOACC_PARALLEL))
+ {
+ gcc_assert ((c = omp_find_clause (clauses, OMP_CLAUSE_MAP))
+ && GOMP_MAP_NONCONTIG_ARRAY_P (OMP_CLAUSE_MAP_KIND (c)));
+ //if (tagging)
+ //args.safe_push (oacc_launch_pack (GOMP_LAUNCH_NONCONTIG_ARRAYS,
+ // NULL_TREE, 0));
+ for (int i = 3; i < TREE_VEC_LENGTH (t); i++)
+ args.safe_push (TREE_VEC_ELT (t, i));
+ }
g = gimple_build_call_vec (builtin_decl_explicit (start_ix), args);
gimple_set_location (g, gimple_location (entry_stmt));
Index: gcc/omp-low.c
===================================================================
--- gcc/omp-low.c (revision 278656)
+++ gcc/omp-low.c (working copy)
@@ -903,6 +903,141 @@ omp_copy_decl (tree var, copy_body_data *cb)
return error_mark_node;
}
+/* Helper function for create_noncontig_array_descr_type(), to append a new field
+ to a record type. */
+
+static void
+append_field_to_record_type (tree record_type, tree fld_ident, tree fld_type)
+{
+ tree *p, fld = build_decl (UNKNOWN_LOCATION, FIELD_DECL, fld_ident, fld_type);
+ DECL_CONTEXT (fld) = record_type;
+
+ for (p = &TYPE_FIELDS (record_type); *p; p = &DECL_CHAIN (*p))
+ ;
+ *p = fld;
+}
+
+/* Create type for non-contiguous array descriptor. Returns created type, and
+ returns the number of dimensions in *DIM_NUM. */
+
+static tree
+create_noncontig_array_descr_type (tree decl, tree dims, int *dim_num)
+{
+ int n = 0;
+ tree array_descr_type, name, x;
+ gcc_assert (TREE_CODE (dims) == TREE_LIST);
+
+ array_descr_type = lang_hooks.types.make_type (RECORD_TYPE);
+ name = create_tmp_var_name (".omp_noncontig_array_descr_type");
+ name = build_decl (UNKNOWN_LOCATION, TYPE_DECL, name, array_descr_type);
+ DECL_ARTIFICIAL (name) = 1;
+ DECL_NAMELESS (name) = 1;
+ TYPE_NAME (array_descr_type) = name;
+ TYPE_ARTIFICIAL (array_descr_type) = 1;
+
+ /* Main starting pointer/array. */
+ /*
+ tree main_var_type = TREE_TYPE (decl);
+ if (TREE_CODE (main_var_type) == REFERENCE_TYPE)
+ main_var_type = TREE_TYPE (main_var_type);
+ append_field_to_record_type (array_descr_type, DECL_NAME (decl),
+ (TREE_CODE (TREE_TYPE (decl)) == POINTER_TYPE
+ ? main_var_type
+ : build_pointer_type (main_var_type)));
+ */
+
+ /* Number of dimensions. */
+ append_field_to_record_type (array_descr_type, get_identifier ("__dim_num"),
+ sizetype);
+
+ for (x = dims; x; x = TREE_CHAIN (x), n++)
+ {
+ char *fldname;
+ /* One for the start index. */
+ ASM_FORMAT_PRIVATE_NAME (fldname, "__dim_base", n);
+ append_field_to_record_type (array_descr_type, get_identifier (fldname),
+ sizetype);
+ /* One for the length. */
+ ASM_FORMAT_PRIVATE_NAME (fldname, "__dim_length", n);
+ append_field_to_record_type (array_descr_type, get_identifier (fldname),
+ sizetype);
+ /* One for the element size. */
+ ASM_FORMAT_PRIVATE_NAME (fldname, "__dim_elem_size", n);
+ append_field_to_record_type (array_descr_type, get_identifier (fldname),
+ sizetype);
+ /* One for is_array flag. */
+ ASM_FORMAT_PRIVATE_NAME (fldname, "__dim_is_array", n);
+ append_field_to_record_type (array_descr_type, get_identifier (fldname),
+ sizetype);
+ }
+
+ layout_type (array_descr_type);
+ *dim_num = n;
+ return array_descr_type;
+}
+
+/* Generate code sequence for initializing non-contiguous array descriptor. */
+
+static void
+create_noncontig_array_descr_init_code (tree array_descr, tree array_var,
+ tree dimensions, int dim_num,
+ gimple_seq *ilist)
+{
+ tree fld, fldref;
+ tree array_descr_type = TREE_TYPE (array_descr);
+ tree dim_type = TREE_TYPE (array_var);
+
+ fld = TYPE_FIELDS (array_descr_type);
+ /*
+ fldref = omp_build_component_ref (array_descr, fld);
+ gimplify_assign (fldref, (TREE_CODE (dim_type) == ARRAY_TYPE
+ ? build_fold_addr_expr (array_var) : array_var),
+ ilist);
+ */
+ if (TREE_CODE (dim_type) == REFERENCE_TYPE)
+ dim_type = TREE_TYPE (dim_type);
+
+ //fld = TREE_CHAIN (fld);
+ fldref = omp_build_component_ref (array_descr, fld);
+ gimplify_assign (fldref, build_int_cst (sizetype, dim_num), ilist);
+
+ while (dimensions)
+ {
+ tree dim_base = fold_convert (sizetype, TREE_PURPOSE (dimensions));
+ tree dim_length = fold_convert (sizetype, TREE_VALUE (dimensions));
+ tree dim_elem_size = TYPE_SIZE_UNIT (TREE_TYPE (dim_type));
+ tree dim_is_array = (TREE_CODE (dim_type) == ARRAY_TYPE
+ ? integer_one_node : integer_zero_node);
+ /* Set base. */
+ fld = TREE_CHAIN (fld);
+ fldref = omp_build_component_ref (array_descr, fld);
+ dim_base = fold_build2 (MULT_EXPR, sizetype, dim_base, dim_elem_size);
+ gimplify_assign (fldref, dim_base, ilist);
+
+ /* Set length. */
+ fld = TREE_CHAIN (fld);
+ fldref = omp_build_component_ref (array_descr, fld);
+ dim_length = fold_build2 (MULT_EXPR, sizetype, dim_length, dim_elem_size);
+ gimplify_assign (fldref, dim_length, ilist);
+
+ /* Set elem_size. */
+ fld = TREE_CHAIN (fld);
+ fldref = omp_build_component_ref (array_descr, fld);
+ dim_elem_size = fold_convert (sizetype, dim_elem_size);
+ gimplify_assign (fldref, dim_elem_size, ilist);
+
+ /* Set is_array flag. */
+ fld = TREE_CHAIN (fld);
+ fldref = omp_build_component_ref (array_descr, fld);
+ dim_is_array = fold_convert (sizetype, dim_is_array);
+ gimplify_assign (fldref, dim_is_array, ilist);
+
+ dimensions = TREE_CHAIN (dimensions);
+ dim_type = TREE_TYPE (dim_type);
+ }
+ gcc_assert (TREE_CHAIN (fld) == NULL_TREE);
+}
+
/* Create a new context, with OUTER_CTX being the surrounding context. */
static omp_context *
@@ -1385,6 +1520,38 @@ scan_sharing_clauses (tree clauses, omp_context *c
install_var_local (decl, ctx);
break;
}
+
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && GOMP_MAP_NONCONTIG_ARRAY_P (OMP_CLAUSE_MAP_KIND (c)))
+ {
+ tree array_decl = OMP_CLAUSE_DECL (c);
+ tree array_type = TREE_TYPE (array_decl);
+ bool by_ref = (TREE_CODE (array_type) == ARRAY_TYPE
+ ? true : false);
+
+ /* Checking code to ensure we only have arrays at top dimension.
+ This limitation might be lifted in the future. See PR76639. */
+ if (TREE_CODE (array_type) == REFERENCE_TYPE)
+ array_type = TREE_TYPE (array_type);
+ tree t = array_type, prev_t = NULL_TREE;
+ while (t)
+ {
+ if (TREE_CODE (t) == ARRAY_TYPE && prev_t)
+ {
+ error_at (gimple_location (ctx->stmt), "array types are"
+ " only allowed at outermost dimension of"
+ " non-contiguous array");
+ break;
+ }
+ prev_t = t;
+ t = TREE_TYPE (t);
+ }
+
+ install_var_field (array_decl, by_ref, 3, ctx);
+ install_var_local (array_decl, ctx);
+ break;
+ }
+
if (DECL_P (decl))
{
if (DECL_SIZE (decl)
@@ -2697,6 +2864,50 @@ scan_omp_single (gomp_single *stmt, omp_context *o
layout_type (ctx->record_type);
}
+/* Reorder clauses so that non-contiguous array map clauses are placed at the very
+ front of the chain. */
+
+static void
+reorder_noncontig_array_clauses (tree *clauses_ptr)
+{
+ tree c, clauses = *clauses_ptr;
+ tree prev_clause = NULL_TREE, next_clause;
+ tree array_clauses = NULL_TREE, array_clauses_tail = NULL_TREE;
+
+ for (c = clauses; c; c = next_clause)
+ {
+ next_clause = OMP_CLAUSE_CHAIN (c);
+
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && GOMP_MAP_NONCONTIG_ARRAY_P (OMP_CLAUSE_MAP_KIND (c)))
+ {
+ /* Unchain c from clauses. */
+ if (c == clauses)
+ clauses = next_clause;
+
+ /* Link on to array_clauses. */
+ if (array_clauses_tail)
+ OMP_CLAUSE_CHAIN (array_clauses_tail) = c;
+ else
+ array_clauses = c;
+ array_clauses_tail = c;
+
+ if (prev_clause)
+ OMP_CLAUSE_CHAIN (prev_clause) = next_clause;
+ continue;
+ }
+
+ prev_clause = c;
+ }
+
+ /* Place non-contiguous array clauses at the start of the clause list. */
+ if (array_clauses)
+ {
+ OMP_CLAUSE_CHAIN (array_clauses_tail) = clauses;
+ *clauses_ptr = array_clauses;
+ }
+}
+
/* Scan a GIMPLE_OMP_TARGET. */
static void
@@ -2705,7 +2916,6 @@ scan_omp_target (gomp_target *stmt, omp_context *o
omp_context *ctx;
tree name;
bool offloaded = is_gimple_omp_offloaded (stmt);
- tree clauses = gimple_omp_target_clauses (stmt);
ctx = new_omp_context (stmt, outer_ctx);
ctx->field_map = splay_tree_new (splay_tree_compare_pointers, 0, 0);
@@ -2724,6 +2934,14 @@ scan_omp_target (gomp_target *stmt, omp_context *o
gimple_omp_target_set_child_fn (stmt, ctx->cb.dst_fn);
}
+ /* If is OpenACC construct, put non-contiguous array clauses (if any)
+ in front of clause chain. The runtime can then test the first to see
+ if the additional map processing for them is required. */
+ if (is_gimple_omp_oacc (stmt))
+ reorder_noncontig_array_clauses (gimple_omp_target_clauses_ptr (stmt));
+
+ tree clauses = gimple_omp_target_clauses (stmt);
+
scan_sharing_clauses (clauses, ctx);
scan_omp (gimple_omp_body_ptr (stmt), ctx);
@@ -11438,6 +11656,15 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp
case GOMP_MAP_FORCE_PRESENT:
case GOMP_MAP_FORCE_DEVICEPTR:
case GOMP_MAP_DEVICE_RESIDENT:
+ case GOMP_MAP_NONCONTIG_ARRAY_TO:
+ case GOMP_MAP_NONCONTIG_ARRAY_FROM:
+ case GOMP_MAP_NONCONTIG_ARRAY_TOFROM:
+ case GOMP_MAP_NONCONTIG_ARRAY_FORCE_TO:
+ case GOMP_MAP_NONCONTIG_ARRAY_FORCE_FROM:
+ case GOMP_MAP_NONCONTIG_ARRAY_FORCE_TOFROM:
+ case GOMP_MAP_NONCONTIG_ARRAY_ALLOC:
+ case GOMP_MAP_NONCONTIG_ARRAY_FORCE_ALLOC:
+ case GOMP_MAP_NONCONTIG_ARRAY_FORCE_PRESENT:
case GOMP_MAP_LINK:
gcc_assert (is_gimple_omp_oacc (stmt));
break;
@@ -11500,7 +11727,14 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp
if (offloaded && !(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_IN_REDUCTION (c)))
{
- x = build_receiver_ref (var, true, ctx);
+ tree var_type = TREE_TYPE (var);
+ bool rcv_by_ref =
+ (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && GOMP_MAP_NONCONTIG_ARRAY_P (OMP_CLAUSE_MAP_KIND (c))
+ && TREE_CODE (var_type) != ARRAY_TYPE
+ ? false : true);
+
+ x = build_receiver_ref (var, rcv_by_ref, ctx);
tree new_var = lookup_decl (var, ctx);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
@@ -11674,6 +11908,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp
vec_alloc (vkind, map_cnt);
unsigned int map_idx = 0;
+ vec<tree> nca_descrs = vNULL;
+
for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
switch (OMP_CLAUSE_CODE (c))
{
@@ -11750,6 +11986,28 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp
avar = build_fold_addr_expr (avar);
gimplify_assign (x, avar, &ilist);
}
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && GOMP_MAP_NONCONTIG_ARRAY_P (OMP_CLAUSE_MAP_KIND (c)))
+ {
+ int dim_num;
+ tree dimensions = OMP_CLAUSE_SIZE (c);
+
+ tree array_descr_type =
+ create_noncontig_array_descr_type (OMP_CLAUSE_DECL (c),
+ dimensions, &dim_num);
+ tree array_descr =
+ create_tmp_var_raw (array_descr_type,
+ ".omp_noncontig_array_descr");
+ gimple_add_tmp_var (array_descr);
+
+ create_noncontig_array_descr_init_code
+ (array_descr, ovar, dimensions, dim_num, &ilist);
+ nca_descrs.safe_push (build_fold_addr_expr (array_descr));
+
+ gimplify_assign (x, (TREE_CODE (TREE_TYPE (ovar)) == ARRAY_TYPE
+ ? build_fold_addr_expr (ovar) : ovar),
+ &ilist);
+ }
else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
{
gcc_assert (is_gimple_omp_oacc (ctx->stmt));
@@ -11821,6 +12079,9 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp
s = TREE_TYPE (s);
s = TYPE_SIZE_UNIT (s);
}
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && GOMP_MAP_NONCONTIG_ARRAY_P (OMP_CLAUSE_MAP_KIND (c)))
+ s = NULL_TREE;
else
s = OMP_CLAUSE_SIZE (c);
if (s == NULL_TREE)
@@ -12072,6 +12333,19 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp
gcc_assert (map_idx == map_cnt);
+ unsigned nca_num = nca_descrs.length ();
+ if (nca_num > 0)
+ {
+ tree nca, t = gimple_omp_target_data_arg (stmt);
+ int i, oldlen = TREE_VEC_LENGTH (t);
+ tree nt = make_tree_vec (oldlen + nca_num);
+ for (i = 0; i < oldlen; i++)
+ TREE_VEC_ELT (nt, i) = TREE_VEC_ELT (t, i);
+ for (i = 0; nca_descrs.iterate (i, &nca); i++)
+ TREE_VEC_ELT (nt, oldlen + i) = nca;
+ gimple_omp_target_set_data_arg (stmt, nt);
+ }
+
DECL_INITIAL (TREE_VEC_ELT (t, 1))
= build_constructor (TREE_TYPE (TREE_VEC_ELT (t, 1)), vsize);
DECL_INITIAL (TREE_VEC_ELT (t, 2))
Index: gcc/testsuite/c-c++-common/goacc/noncontig_array-1.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/noncontig_array-1.c (nonexistent)
+++ gcc/testsuite/c-c++-common/goacc/noncontig_array-1.c (working copy)
@@ -0,0 +1,25 @@
+/* { dg-do compile } */
+
+void foo (void)
+{
+ int array_of_array[10][10];
+ int **ptr_to_ptr;
+ int *array_of_ptr[10];
+ int (*ptr_to_array)[10];
+
+ #pragma acc parallel copy (array_of_array[2:4][0:10])
+ array_of_array[5][5] = 1;
+
+ #pragma acc parallel copy (ptr_to_ptr[2:4][1:7])
+ ptr_to_ptr[5][5] = 1;
+
+ #pragma acc parallel copy (array_of_ptr[2:4][1:7])
+ array_of_ptr[5][5] = 1;
+
+ #pragma acc parallel copy (ptr_to_array[2:4][1:7]) /* { dg-error "array section is not contiguous in 'map' clause" } */
+ ptr_to_array[5][5] = 1;
+}
+/* { dg-final { scan-tree-dump-times {#pragma omp target oacc_parallel map\(tofrom:array_of_array} 1 gimple } } */
+/* { dg-final { scan-tree-dump-times {#pragma omp target oacc_parallel map\(tofrom,noncontig_array:ptr_to_ptr \[dimensions: 2 4, 1 7\]} 1 gimple } } */
+/* { dg-final { scan-tree-dump-times {#pragma omp target oacc_parallel map\(tofrom,noncontig_array:array_of_ptr \[dimensions: 2 4, 1 7\]} 1 gimple } } */
+/* { dg-final { scan-tree-dump-times {#pragma omp target oacc_parallel map\(tofrom,noncontig_array:ptr_to_array \[dimensions: 2 4, 1 7\]} 1 gimple { xfail *-*-* } } } */
Index: gcc/tree-pretty-print.c
===================================================================
--- gcc/tree-pretty-print.c (revision 278656)
+++ gcc/tree-pretty-print.c (working copy)
@@ -849,6 +849,33 @@ dump_omp_clause (pretty_printer *pp, tree clause,
case GOMP_MAP_LINK:
pp_string (pp, "link");
break;
+ case GOMP_MAP_NONCONTIG_ARRAY_TO:
+ pp_string (pp, "to,noncontig_array");
+ break;
+ case GOMP_MAP_NONCONTIG_ARRAY_FROM:
+ pp_string (pp, "from,noncontig_array");
+ break;
+ case GOMP_MAP_NONCONTIG_ARRAY_TOFROM:
+ pp_string (pp, "tofrom,noncontig_array");
+ break;
+ case GOMP_MAP_NONCONTIG_ARRAY_FORCE_TO:
+ pp_string (pp, "force_to,noncontig_array");
+ break;
+ case GOMP_MAP_NONCONTIG_ARRAY_FORCE_FROM:
+ pp_string (pp, "force_from,noncontig_array");
+ break;
+ case GOMP_MAP_NONCONTIG_ARRAY_FORCE_TOFROM:
+ pp_string (pp, "force_tofrom,noncontig_array");
+ break;
+ case GOMP_MAP_NONCONTIG_ARRAY_ALLOC:
+ pp_string (pp, "alloc,noncontig_array");
+ break;
+ case GOMP_MAP_NONCONTIG_ARRAY_FORCE_ALLOC:
+ pp_string (pp, "force_alloc,noncontig_array");
+ break;
+ case GOMP_MAP_NONCONTIG_ARRAY_FORCE_PRESENT:
+ pp_string (pp, "force_present,noncontig_array");
+ break;
default:
gcc_unreachable ();
}
@@ -859,8 +886,15 @@ dump_omp_clause (pretty_printer *pp, tree clause,
if (OMP_CLAUSE_SIZE (clause))
{
switch (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP
- ? OMP_CLAUSE_MAP_KIND (clause) : GOMP_MAP_TO)
+ ? (GOMP_MAP_NONCONTIG_ARRAY_P (OMP_CLAUSE_MAP_KIND (clause))
+ ? GOMP_MAP_NONCONTIG_ARRAY
+ : OMP_CLAUSE_MAP_KIND (clause))
+ : GOMP_MAP_TO)
{
+ case GOMP_MAP_NONCONTIG_ARRAY:
+ gcc_assert (TREE_CODE (OMP_CLAUSE_SIZE (clause)) == TREE_LIST);
+ pp_string (pp, " [dimensions: ");
+ break;
case GOMP_MAP_POINTER:
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
Index: libgomp/libgomp.h
===================================================================
--- libgomp/libgomp.h (revision 278656)
+++ libgomp/libgomp.h (working copy)
@@ -1167,6 +1167,10 @@ extern struct target_mem_desc *gomp_map_vars_async
size_t, void **, void **,
size_t *, void *, bool,
enum gomp_map_vars_kind);
+extern struct target_mem_desc *gomp_map_vars_openacc (struct gomp_device_descr *,
+ struct goacc_asyncqueue *,
+ size_t, void **, size_t *,
+ unsigned short *, void *);
extern void gomp_unmap_tgt (struct target_mem_desc *);
extern void gomp_unmap_vars (struct target_mem_desc *, bool);
extern void gomp_unmap_vars_async (struct target_mem_desc *, bool,
Index: libgomp/libgomp_g.h
===================================================================
--- libgomp/libgomp_g.h (revision 278656)
+++ libgomp/libgomp_g.h (working copy)
@@ -364,7 +364,7 @@ extern void GOACC_parallel_keyed (int, void (*) (v
extern void GOACC_parallel (int, void (*) (void *), size_t, void **, size_t *,
unsigned short *, int, int, int, int, int, ...);
extern void GOACC_data_start (int, size_t, void **, size_t *,
- unsigned short *);
+ unsigned short *, ...);
extern void GOACC_data_end (void);
extern void GOACC_enter_exit_data (int, size_t, void **,
size_t *, unsigned short *, int, int, ...);
Index: libgomp/oacc-int.h
===================================================================
--- libgomp/oacc-int.h (revision 278656)
+++ libgomp/oacc-int.h (working copy)
@@ -164,6 +164,57 @@ bool _goacc_profiling_setup_p (struct goacc_thread
void goacc_profiling_dispatch (acc_prof_info *, acc_event_info *,
acc_api_info *);
+/* Definitions for data structures describing OpenACC non-contiguous arrays
+ (Note: interfaces with compiler)
+
+ The compiler generates a descriptor for each such array, places the
+ descriptor on stack, and passes the address of the descriptor to the libgomp
+ runtime as a normal map argument. The runtime then processes the array
+ data structure setup, and replaces the argument with the new actual
+ array address for the child function.
+
+ Care must be taken such that the struct field and layout assumptions
+ of struct goacc_ncarray_dim, goacc_ncarray_descr_type inside the compiler
+ be consistant with the below declarations. */
+
+struct goacc_ncarray_dim {
+ size_t base;
+ size_t length;
+ size_t elem_size;
+ size_t is_array;
+};
+
+struct goacc_ncarray_descr_type
+{
+ size_t ndims;
+ struct goacc_ncarray_dim dims[];
+};
+
+/* Internal non-contiguous array info struct, used only here inside the runtime. */
+
+struct goacc_ncarray
+{
+ struct goacc_ncarray_descr_type *descr;
+ void *ptr;
+ size_t map_index;
+ size_t ptrblock_size;
+ void **data_rows;
+ void **tgt_data_rows;
+ size_t data_row_num;
+ size_t data_row_size;
+};
+
+struct goacc_ncarray_info
+{
+ size_t num_data_rows, num_ncarray;
+ void **data_rows;
+ void **tgt_data_rows;
+ struct goacc_ncarray ncarray[];
+};
+
+extern void *goacc_noncontig_array_create_ptrblock (struct goacc_ncarray *, void *);
+
+
#ifdef HAVE_ATTRIBUTE_VISIBILITY
# pragma GCC visibility pop
#endif
Index: libgomp/oacc-parallel.c
===================================================================
--- libgomp/oacc-parallel.c (revision 278656)
+++ libgomp/oacc-parallel.c (working copy)
@@ -36,8 +36,8 @@
#include <string.h>
#include <stdarg.h>
#include <assert.h>
+#include <stdio.h>
-
/* In the ABI, the GOACC_FLAGs are encoded as an inverted bitmask, so that we
continue to support the following two legacy values. */
_Static_assert (GOACC_FLAGS_UNMARSHAL (GOMP_DEVICE_ICV) == 0,
@@ -113,7 +113,174 @@ handle_ftn_pointers (size_t mapnum, void **hostadd
static void goacc_wait (int async, int num_waits, va_list *ap);
+static size_t
+goacc_noncontig_array_count_rows (struct goacc_ncarray_descr_type *descr)
+{
+ size_t nrows = 1;
+ for (size_t d = 0; d < descr->ndims - 1; d++)
+ nrows *= descr->dims[d].length / sizeof (void *);
+ return nrows;
+}
+static void
+goacc_noncontig_array_compute_sizes (struct goacc_ncarray *nca)
+{
+ size_t d, n = 1;
+ struct goacc_ncarray_descr_type *descr = nca->descr;
+
+ nca->ptrblock_size = 0;
+ for (d = 0; d < descr->ndims - 1; d++)
+ {
+ size_t dim_count = descr->dims[d].length / descr->dims[d].elem_size;
+ size_t dim_ptrblock_size = (descr->dims[d + 1].is_array
+ ? 0 : descr->dims[d].length * n);
+ nca->ptrblock_size += dim_ptrblock_size;
+ n *= dim_count;
+ }
+ nca->data_row_num = n;
+ nca->data_row_size = descr->dims[d].length;
+
+}
+
+static void
+goacc_noncontig_array_fill_rows_1 (struct goacc_ncarray_descr_type *descr, void *nca,
+ size_t d, void ***row_ptr, size_t *count)
+{
+ if (d < descr->ndims - 1)
+ {
+ size_t elsize = descr->dims[d].elem_size;
+ size_t n = descr->dims[d].length / elsize;
+ void *p = nca + descr->dims[d].base;
+ for (size_t i = 0; i < n; i++)
+ {
+ void *ptr = p + i * elsize;
+ /* Deref if next dimension is not array. */
+ if (!descr->dims[d + 1].is_array)
+ ptr = *((void **) ptr);
+ goacc_noncontig_array_fill_rows_1 (descr, ptr, d + 1, row_ptr, count);
+ }
+ }
+ else
+ {
+ **row_ptr = nca + descr->dims[d].base;
+ *row_ptr += 1;
+ *count += 1;
+ }
+}
+
+static size_t
+goacc_noncontig_array_fill_rows (struct goacc_ncarray *nca)
+{
+ size_t count = 0;
+ void **p = nca->data_rows;
+ goacc_noncontig_array_fill_rows_1 (nca->descr, nca->ptr, 0, &p, &count);
+ return count;
+}
+
+static struct goacc_ncarray_info *
+goacc_process_noncontiguous_arrays (size_t mapnum, void **hostaddrs,
+ unsigned short *kinds, va_list* ap)
+{
+ size_t i, nr, num_data_rows = 0, num_ncarray = 0, curr_row_start = 0;
+ struct goacc_ncarray_descr_type *descr;
+
+ /* We need to go over *ap twice, so preserve *ap state here. */
+ va_list itr;
+ va_copy (itr, *ap);
+ for (i = 0; i < mapnum; i++)
+ if (GOMP_MAP_NONCONTIG_ARRAY_P (kinds[i] & 0xff))
+ {
+ descr = va_arg (itr, struct goacc_ncarray_descr_type *);
+ num_data_rows += goacc_noncontig_array_count_rows (descr);
+ num_ncarray += 1;
+ }
+ else
+ break;
+
+ /* Allocate the entire info struct, array entries, and row pointer
+ arrays in one large block. */
+ struct goacc_ncarray_info *nca_info
+ = gomp_malloc (sizeof (struct goacc_ncarray_info)
+ + sizeof (struct goacc_ncarray) * num_ncarray
+ + sizeof (void *) * num_data_rows * 2);
+ nca_info->num_data_rows = num_data_rows;
+ nca_info->num_ncarray = num_ncarray;
+ nca_info->data_rows = (void **) (nca_info->ncarray + num_ncarray);
+ nca_info->tgt_data_rows = nca_info->data_rows + num_data_rows;
+
+ struct goacc_ncarray *curr_ncarray = nca_info->ncarray;
+ for (i = 0; i < mapnum; i++)
+ if (GOMP_MAP_NONCONTIG_ARRAY_P (kinds[i] & 0xff))
+ {
+ descr = va_arg (*ap, struct goacc_ncarray_descr_type *);
+ curr_ncarray->descr = descr;
+ curr_ncarray->ptr = hostaddrs[i];
+ curr_ncarray->map_index = i;
+
+ goacc_noncontig_array_compute_sizes (curr_ncarray);
+
+ curr_ncarray->data_rows = nca_info->data_rows + curr_row_start;
+ curr_ncarray->tgt_data_rows = nca_info->tgt_data_rows + curr_row_start;
+
+ nr = goacc_noncontig_array_fill_rows (curr_ncarray);
+ assert (nr == curr_ncarray->data_row_num);
+ curr_row_start += nr;
+ curr_ncarray += 1;
+ }
+ else
+ break;
+
+ return nca_info;
+}
+
+void *
+goacc_noncontig_array_create_ptrblock (struct goacc_ncarray *nca,
+ void *tgt_ptrblock_addr)
+{
+ struct goacc_ncarray_descr_type *descr = nca->descr;
+ void **tgt_data_rows = nca->tgt_data_rows;
+ void *ptrblock = gomp_malloc (nca->ptrblock_size);
+ void **curr_dim_ptrblock = (void **) ptrblock;
+ size_t n = 1;
+
+ for (size_t d = 0; d < descr->ndims - 1; d++)
+ {
+ int curr_dim_len = descr->dims[d].length;
+ int next_dim_len = descr->dims[d + 1].length;
+ int curr_dim_num = curr_dim_len / sizeof (void *);
+ size_t next_dim_bias = descr->dims[d + 1].base;
+
+ void *next_dim_ptrblock
+ = (void *)(curr_dim_ptrblock + n * curr_dim_num);
+
+ for (int b = 0; b < n; b++)
+ for (int i = 0; i < curr_dim_num; i++)
+ {
+ if (d < descr->ndims - 2)
+ {
+ void *ptr = (next_dim_ptrblock
+ + b * curr_dim_num * next_dim_len
+ + i * next_dim_len);
+ void *tgt_ptr = (tgt_ptrblock_addr
+ + (ptr - ptrblock) - next_dim_bias);
+ curr_dim_ptrblock[b * curr_dim_num + i] = tgt_ptr;
+ }
+ else
+ {
+ curr_dim_ptrblock[b * curr_dim_num + i]
+ = tgt_data_rows[b * curr_dim_num + i] - next_dim_bias;
+ }
+ void *addr = &curr_dim_ptrblock[b * curr_dim_num + i];
+ assert (ptrblock <= addr && addr < ptrblock + nca->ptrblock_size);
+ }
+
+ n *= curr_dim_num;
+ curr_dim_ptrblock = next_dim_ptrblock;
+ }
+ assert (n == nca->data_row_num);
+ return ptrblock;
+}
+
/* Launch a possibly offloaded function with FLAGS. FN is the host fn
address. MAPNUM, HOSTADDRS, SIZES & KINDS describe the memory
blocks to be copied to/from the device. Varadic arguments are
@@ -138,6 +305,7 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (voi
int async = GOMP_ASYNC_SYNC;
unsigned dims[GOMP_DIM_MAX];
unsigned tag;
+ struct goacc_ncarray_info *nca_info = NULL;
#ifdef HAVE_INTTYPES_H
gomp_debug (0, "%s: mapnum=%"PRIu64", hostaddrs=%p, size=%p, kinds=%p\n",
@@ -270,11 +438,20 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (voi
break;
}
+ /*case GOMP_LAUNCH_NONCONTIG_ARRAYS:
+ nca_info = goacc_process_noncontiguous_arrays (mapnum, hostaddrs,
+ kinds, &ap);
+ break;*/
+
default:
gomp_fatal ("unrecognized offload code '%d',"
" libgomp is too old", GOMP_LAUNCH_CODE (tag));
}
}
+
+ if (mapnum > 0 && GOMP_MAP_NONCONTIG_ARRAY_P (kinds[0] & 0xff))
+ nca_info = goacc_process_noncontiguous_arrays (mapnum, hostaddrs, kinds, &ap);
+
va_end (ap);
if (!(acc_dev->capabilities & GOMP_OFFLOAD_CAP_NATIVE_EXEC))
@@ -311,8 +488,10 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (voi
goacc_aq aq = get_goacc_asyncqueue (async);
- tgt = gomp_map_vars_async (acc_dev, aq, mapnum, hostaddrs, NULL, sizes, kinds,
- true, GOMP_MAP_VARS_OPENACC);
+ tgt = gomp_map_vars_openacc (acc_dev, aq, mapnum, hostaddrs, sizes, kinds,
+ nca_info);
+ free (nca_info);
+
if (profiling_p)
{
prof_info.event_type = acc_ev_enter_data_end;
@@ -390,7 +569,7 @@ GOACC_parallel (int flags_m, void (*fn) (void *),
void
GOACC_data_start (int flags_m, size_t mapnum,
- void **hostaddrs, size_t *sizes, unsigned short *kinds)
+ void **hostaddrs, size_t *sizes, unsigned short *kinds, ...)
{
int flags = GOACC_FLAGS_UNMARSHAL (flags_m);
@@ -480,8 +659,7 @@ GOACC_data_start (int flags_m, size_t mapnum,
{
prof_info.device_type = acc_device_host;
api_info.device_type = prof_info.device_type;
- tgt = gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, true,
- GOMP_MAP_VARS_OPENACC);
+ tgt = gomp_map_vars_openacc (NULL, NULL, 0, NULL, NULL, NULL, NULL);
tgt->prev = thr->mapped_data;
thr->mapped_data = tgt;
@@ -488,9 +666,19 @@ GOACC_data_start (int flags_m, size_t mapnum,
goto out_prof;
}
+ struct goacc_ncarray_info *nca_info = NULL;
+ if (mapnum > 0 && GOMP_MAP_NONCONTIG_ARRAY_P (kinds[0] & 0xff))
+ {
+ va_list ap;
+ va_start (ap, kinds);
+ nca_info = goacc_process_noncontiguous_arrays (mapnum, hostaddrs, kinds, &ap);
+ va_end (ap);
+ }
+
gomp_debug (0, " %s: prepare mappings\n", __FUNCTION__);
- tgt = gomp_map_vars (acc_dev, mapnum, hostaddrs, NULL, sizes, kinds, true,
- GOMP_MAP_VARS_OPENACC);
+ tgt = gomp_map_vars_openacc (acc_dev, NULL, mapnum, hostaddrs, sizes, kinds,
+ nca_info);
+ free (nca_info);
gomp_debug (0, " %s: mappings prepared\n", __FUNCTION__);
tgt->prev = thr->mapped_data;
thr->mapped_data = tgt;
@@ -701,7 +889,6 @@ GOACC_enter_exit_data (int flags_m, size_t mapnum,
if (num_waits)
{
va_list ap;
-
va_start (ap, num_waits);
goacc_wait (async, num_waits, &ap);
va_end (ap);
Index: libgomp/target.c
===================================================================
--- libgomp/target.c (revision 278656)
+++ libgomp/target.c (working copy)
@@ -524,10 +524,11 @@ static inline __attribute__((always_inline)) struc
gomp_map_vars_internal (struct gomp_device_descr *devicep,
struct goacc_asyncqueue *aq, size_t mapnum,
void **hostaddrs, void **devaddrs, size_t *sizes,
- void *kinds, bool short_mapkind,
- enum gomp_map_vars_kind pragma_kind)
+ void *kinds, struct goacc_ncarray_info *nca_info,
+ bool short_mapkind, enum gomp_map_vars_kind pragma_kind)
{
size_t i, tgt_align, tgt_size, not_found_cnt = 0;
+ size_t nca_data_row_num = (nca_info ? nca_info->num_data_rows : 0);
bool has_firstprivate = false;
const int rshift = short_mapkind ? 8 : 3;
const int typemask = short_mapkind ? 0xff : 0x7;
@@ -534,8 +535,9 @@ gomp_map_vars_internal (struct gomp_device_descr *
struct splay_tree_s *mem_map = &devicep->mem_map;
struct splay_tree_key_s cur_node;
struct target_mem_desc *tgt
- = gomp_malloc (sizeof (*tgt) + sizeof (tgt->list[0]) * mapnum);
- tgt->list_count = mapnum;
+ = gomp_malloc (sizeof (*tgt)
+ + sizeof (tgt->list[0]) * (mapnum + nca_data_row_num));
+ tgt->list_count = mapnum + nca_data_row_num;
tgt->refcount = pragma_kind == GOMP_MAP_VARS_ENTER_DATA ? 0 : 1;
tgt->device_descr = devicep;
struct gomp_coalesce_buf cbuf, *cbufp = NULL;
@@ -667,6 +669,27 @@ gomp_map_vars_internal (struct gomp_device_descr *
has_firstprivate = true;
continue;
}
+ else if (GOMP_MAP_NONCONTIG_ARRAY_P (kind & typemask))
+ {
+ /* Ignore non-contiguous arrays for now, we process them together
+ later. */
+ tgt->list[i].key = NULL;
+ tgt->list[i].offset = 0;
+ not_found_cnt++;
+
+ /* The map for the non-contiguous array itself is never copied from
+ during unmapping, its the data rows that count. Set copy-from
+ flags to false here. */
+ tgt->list[i].copy_from = false;
+ tgt->list[i].always_copy_from = false;
+
+ size_t align = (size_t) 1 << (kind >> rshift);
+ if (tgt_align < align)
+ tgt_align = align;
+
+ continue;
+ }
+
cur_node.host_start = (uintptr_t) hostaddrs[i];
if (!GOMP_MAP_POINTER_P (kind & typemask))
cur_node.host_end = cur_node.host_start + sizes[i];
@@ -735,6 +758,44 @@ gomp_map_vars_internal (struct gomp_device_descr *
}
}
+ /* For non-contiguous arrays. Each data row is one target item, separated
+ from the normal map clause items, hence we order them after mapnum. */
+ if (nca_info)
+ {
+ struct target_var_desc *next_var_desc = &tgt->list[mapnum];
+ for (i = 0; i < nca_info->num_ncarray; i++)
+ {
+ struct goacc_ncarray *nca = &nca_info->ncarray[i];
+ int kind = get_kind (short_mapkind, kinds, nca->map_index);
+ size_t align = (size_t) 1 << (kind >> rshift);
+ tgt_size = (tgt_size + align - 1) & ~(align - 1);
+ tgt_size += nca->ptrblock_size;
+
+ for (size_t j = 0; j < nca->data_row_num; j++)
+ {
+ struct target_var_desc *row_desc = next_var_desc++;
+ void *row = nca->data_rows[j];
+ cur_node.host_start = (uintptr_t) row;
+ cur_node.host_end = cur_node.host_start + nca->data_row_size;
+ splay_tree_key n = splay_tree_lookup (mem_map, &cur_node);
+ if (n)
+ {
+ assert (n->refcount != REFCOUNT_LINK);
+ gomp_map_vars_existing (devicep, aq, n, &cur_node, row_desc,
+ kind & typemask,
+ /* TODO: cbuf? */ NULL);
+ }
+ else
+ {
+ tgt_size = (tgt_size + align - 1) & ~(align - 1);
+ tgt_size += nca->data_row_size;
+ not_found_cnt++;
+ }
+ }
+ }
+ assert (next_var_desc == &tgt->list[mapnum + nca_info->num_data_rows]);
+ }
+
if (devaddrs)
{
if (mapnum != 1)
@@ -895,6 +956,15 @@ gomp_map_vars_internal (struct gomp_device_descr *
default:
break;
}
+
+ if (GOMP_MAP_NONCONTIG_ARRAY_P (kind & typemask))
+ {
+ tgt->list[i].key = &array->key;
+ tgt->list[i].key->tgt = tgt;
+ array++;
+ continue;
+ }
+
splay_tree_key k = &array->key;
k->host_start = (uintptr_t) hostaddrs[i];
if (!GOMP_MAP_POINTER_P (kind & typemask))
@@ -1044,6 +1114,98 @@ gomp_map_vars_internal (struct gomp_device_descr *
array++;
}
}
+
+ /* Processing of non-contiguous array rows. */
+ if (nca_info)
+ {
+ struct target_var_desc *next_var_desc = &tgt->list[mapnum];
+ for (i = 0; i < nca_info->num_ncarray; i++)
+ {
+ struct goacc_ncarray *nca = &nca_info->ncarray[i];
+ int kind = get_kind (short_mapkind, kinds, nca->map_index);
+ size_t align = (size_t) 1 << (kind >> rshift);
+ tgt_size = (tgt_size + align - 1) & ~(align - 1);
+
+ assert (nca->ptr == hostaddrs[nca->map_index]);
+
+ /* For the map of the non-contiguous array itself, adjust so that
+ the passed device address points to the beginning of the
+ ptrblock. Remember to adjust the first-dimension's bias here. */
+ tgt->list[nca->map_index].key->tgt_offset
+ = tgt_size - nca->descr->dims[0].base;
+
+ void *target_ptrblock = (void*) tgt->tgt_start + tgt_size;
+ tgt_size += nca->ptrblock_size;
+
+ /* Add splay key for each data row in current non-contiguous
+ array. */
+ for (size_t j = 0; j < nca->data_row_num; j++)
+ {
+ struct target_var_desc *row_desc = next_var_desc++;
+ void *row = nca->data_rows[j];
+ cur_node.host_start = (uintptr_t) row;
+ cur_node.host_end = cur_node.host_start + nca->data_row_size;
+ splay_tree_key k = splay_tree_lookup (mem_map, &cur_node);
+ if (k)
+ {
+ assert (k->refcount != REFCOUNT_LINK);
+ gomp_map_vars_existing (devicep, aq, k, &cur_node, row_desc,
+ kind & typemask,
+ cbufp);
+ }
+ else
+ {
+ tgt->refcount++;
+ tgt_size = (tgt_size + align - 1) & ~(align - 1);
+
+ k = &array->key;
+ k->host_start = (uintptr_t) row;
+ k->host_end = k->host_start + nca->data_row_size;
+
+ k->tgt = tgt;
+ k->refcount = 1;
+ k->dynamic_refcount = 0;
+ k->link_key = NULL;
+ k->tgt_offset = tgt_size;
+
+ tgt_size += nca->data_row_size;
+
+ row_desc->key = k;
+ row_desc->copy_from
+ = GOMP_MAP_COPY_FROM_P (kind & typemask);
+ row_desc->always_copy_from
+ = GOMP_MAP_COPY_FROM_P (kind & typemask);
+ row_desc->offset = 0;
+ row_desc->length = nca->data_row_size;
+
+ array->left = NULL;
+ array->right = NULL;
+ splay_tree_insert (mem_map, array);
+
+ if (GOMP_MAP_COPY_TO_P (kind & typemask))
+ gomp_copy_host2dev (devicep, aq,
+ (void *) tgt->tgt_start + k->tgt_offset,
+ (void *) k->host_start,
+ nca->data_row_size, cbufp);
+ array++;
+ }
+ nca->tgt_data_rows[j]
+ = (void *) (k->tgt->tgt_start + k->tgt_offset);
+ }
+
+ /* Now we have the target memory allocated, and target offsets of all
+ row blocks assigned and calculated, we can construct the
+ accelerator side ptrblock and copy it in. */
+ if (nca->ptrblock_size)
+ {
+ void *ptrblock = goacc_noncontig_array_create_ptrblock
+ (nca, target_ptrblock);
+ gomp_copy_host2dev (devicep, aq, target_ptrblock, ptrblock,
+ nca->ptrblock_size, cbufp);
+ free (ptrblock);
+ }
+ }
+ }
}
if (pragma_kind == GOMP_MAP_VARS_TARGET)
@@ -1086,12 +1248,25 @@ gomp_map_vars_internal (struct gomp_device_descr *
}
attribute_hidden struct target_mem_desc *
+gomp_map_vars_openacc (struct gomp_device_descr *devicep,
+ struct goacc_asyncqueue *aq, size_t mapnum,
+ void **hostaddrs, size_t *sizes, unsigned short *kinds,
+ void *nca_info)
+{
+ return gomp_map_vars_internal (devicep, aq, mapnum, hostaddrs, NULL,
+ sizes, (void *) kinds,
+ (struct goacc_ncarray_info *) nca_info,
+ true, GOMP_MAP_VARS_OPENACC);
+}
+
+attribute_hidden struct target_mem_desc *
gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
void **hostaddrs, void **devaddrs, size_t *sizes, void *kinds,
bool short_mapkind, enum gomp_map_vars_kind pragma_kind)
{
return gomp_map_vars_internal (devicep, NULL, mapnum, hostaddrs, devaddrs,
- sizes, kinds, short_mapkind, pragma_kind);
+ sizes, kinds, NULL, short_mapkind,
+ pragma_kind);
}
attribute_hidden struct target_mem_desc *
@@ -1102,7 +1277,7 @@ gomp_map_vars_async (struct gomp_device_descr *dev
enum gomp_map_vars_kind pragma_kind)
{
return gomp_map_vars_internal (devicep, aq, mapnum, hostaddrs, devaddrs,
- sizes, kinds, short_mapkind, pragma_kind);
+ sizes, kinds, NULL, short_mapkind, pragma_kind);
}
attribute_hidden void
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-1.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-1.c (nonexistent)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-1.c (working copy)
@@ -0,0 +1,103 @@
+/* { dg-do run } */
+
+#include <stdlib.h>
+#include <assert.h>
+
+#define n 100
+#define m 100
+
+int b[n][m];
+
+void
+test1 (void)
+{
+ int i, j, *a[100];
+
+ /* Array of pointers form test. */
+ for (i = 0; i < n; i++)
+ {
+ a[i] = (int *)malloc (sizeof (int) * m);
+ for (j = 0; j < m; j++)
+ b[i][j] = j - i;
+ }
+
+ #pragma acc parallel loop copyout(a[0:n][0:m]) copyin(b)
+ for (i = 0; i < n; i++)
+ #pragma acc loop
+ for (j = 0; j < m; j++)
+ a[i][j] = b[i][j];
+
+ for (i = 0; i < n; i++)
+ {
+ for (j = 0; j < m; j++)
+ assert (a[i][j] == b[i][j]);
+ /* Clean up. */
+ free (a[i]);
+ }
+}
+
+void
+test2 (void)
+{
+ int i, j, **a = (int **) malloc (sizeof (int *) * n);
+
+ /* Separately allocated blocks. */
+ for (i = 0; i < n; i++)
+ {
+ a[i] = (int *)malloc (sizeof (int) * m);
+ for (j = 0; j < m; j++)
+ b[i][j] = j - i;
+ }
+
+ #pragma acc parallel loop copyout(a[0:n][0:m]) copyin(b)
+ for (i = 0; i < n; i++)
+ #pragma acc loop
+ for (j = 0; j < m; j++)
+ a[i][j] = b[i][j];
+
+ for (i = 0; i < n; i++)
+ {
+ for (j = 0; j < m; j++)
+ assert (a[i][j] == b[i][j]);
+ /* Clean up. */
+ free (a[i]);
+ }
+ free (a);
+}
+
+void
+test3 (void)
+{
+ int i, j, **a = (int **) malloc (sizeof (int *) * n);
+ a[0] = (int *) malloc (sizeof (int) * n * m);
+
+ /* Rows allocated in one contiguous block. */
+ for (i = 0; i < n; i++)
+ {
+ a[i] = *a + i * m;
+ for (j = 0; j < m; j++)
+ b[i][j] = j - i;
+ }
+
+ #pragma acc parallel loop copyout(a[0:n][0:m]) copyin(b)
+ for (i = 0; i < n; i++)
+ #pragma acc loop
+ for (j = 0; j < m; j++)
+ a[i][j] = b[i][j];
+
+ for (i = 0; i < n; i++)
+ for (j = 0; j < m; j++)
+ assert (a[i][j] == b[i][j]);
+
+ free (a[0]);
+ free (a);
+}
+
+int
+main (void)
+{
+ test1 ();
+ test2 ();
+ test3 ();
+ return 0;
+}
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-2.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-2.c (nonexistent)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-2.c (working copy)
@@ -0,0 +1,37 @@
+/* { dg-do run } */
+
+#include <assert.h>
+#include "noncontig_array-utils.h"
+
+int
+main (void)
+{
+ int n = 10;
+ int ***a = (int ***) create_ncarray (sizeof (int), n, 3);
+ int ***b = (int ***) create_ncarray (sizeof (int), n, 3);
+ int ***c = (int ***) create_ncarray (sizeof (int), n, 3);
+
+ for (int i = 0; i < n; i++)
+ for (int j = 0; j < n; j++)
+ for (int k = 0; k < n; k++)
+ {
+ a[i][j][k] = i + j * k + k;
+ b[i][j][k] = j + k * i + i * j;
+ c[i][j][k] = a[i][j][k];
+ }
+
+ #pragma acc parallel copy (a[0:n][0:n][0:n]) copyin (b[0:n][0:n][0:n])
+ {
+ for (int i = 0; i < n; i++)
+ for (int j = 0; j < n; j++)
+ for (int k = 0; k < n; k++)
+ a[i][j][k] += b[k][j][i] + i + j + k;
+ }
+
+ for (int i = 0; i < n; i++)
+ for (int j = 0; j < n; j++)
+ for (int k = 0; k < n; k++)
+ assert (a[i][j][k] == c[i][j][k] + b[k][j][i] + i + j + k);
+
+ return 0;
+}
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-3.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-3.c (nonexistent)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-3.c (working copy)
@@ -0,0 +1,45 @@
+/* { dg-do run } */
+
+#include <assert.h>
+#include "noncontig_array-utils.h"
+
+int main (void)
+{
+ int n = 20, x = 5, y = 12;
+ int *****a = (int *****) create_ncarray (sizeof (int), n, 5);
+
+ int sum1 = 0, sum2 = 0, sum3 = 0;
+
+ for (int i = 0; i < n; i++)
+ for (int j = 0; j < n; j++)
+ for (int k = 0; k < n; k++)
+ for (int l = 0; l < n; l++)
+ for (int m = 0; m < n; m++)
+ {
+ a[i][j][k][l][m] = 1;
+ sum1++;
+ }
+
+ #pragma acc parallel copy (a[x:y][x:y][x:y][x:y][x:y]) copy(sum2)
+ {
+ for (int i = x; i < x + y; i++)
+ for (int j = x; j < x + y; j++)
+ for (int k = x; k < x + y; k++)
+ for (int l = x; l < x + y; l++)
+ for (int m = x; m < x + y; m++)
+ {
+ a[i][j][k][l][m] = 0;
+ sum2++;
+ }
+ }
+
+ for (int i = 0; i < n; i++)
+ for (int j = 0; j < n; j++)
+ for (int k = 0; k < n; k++)
+ for (int l = 0; l < n; l++)
+ for (int m = 0; m < n; m++)
+ sum3 += a[i][j][k][l][m];
+
+ assert (sum1 == sum2 + sum3);
+ return 0;
+}
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-4.c
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-4.c (nonexistent)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-4.c (working copy)
@@ -0,0 +1,36 @@
+/* { dg-do run } */
+
+#include <assert.h>
+#include "noncontig_array-utils.h"
+
+int main (void)
+{
+ int n = 128;
+ double ***a = (double ***) create_ncarray (sizeof (double), n, 3);
+ double ***b = (double ***) create_ncarray (sizeof (double), n, 3);
+
+ for (int i = 0; i < n; i++)
+ for (int j = 0; j < n; j++)
+ for (int k = 0; k < n; k++)
+ a[i][j][k] = i + j + k + i * j * k;
+
+ /* This test exercises async copyout of non-contiguous array rows. */
+ #pragma acc parallel copyin(a[0:n][0:n][0:n]) copyout(b[0:n][0:n][0:n]) async(5)
+ {
+ #pragma acc loop gang
+ for (int i = 0; i < n; i++)
+ #pragma acc loop vector
+ for (int j = 0; j < n; j++)
+ for (int k = 0; k < n; k++)
+ b[i][j][k] = a[i][j][k] * 2.0;
+ }
+
+ #pragma acc wait (5)
+
+ for (int i = 0; i < n; i++)
+ for (int j = 0; j < n; j++)
+ for (int k = 0; k < n; k++)
+ assert (b[i][j][k] == a[i][j][k] * 2.0);
+
+ return 0;
+}
Index: libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-utils.h
===================================================================
--- libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-utils.h (nonexistent)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-utils.h (working copy)
@@ -0,0 +1,44 @@
+#include <stdlib.h>
+#include <string.h>
+#include <assert.h>
+#include <stdint.h>
+
+/* Allocate and create a pointer based NDIMS-dimensional array,
+ each dimension DIMLEN long, with ELSIZE sized data elements. */
+void *
+create_ncarray (size_t elsize, int dimlen, int ndims)
+{
+ size_t blk_size = 0;
+ size_t n = 1;
+
+ for (int i = 0; i < ndims - 1; i++)
+ {
+ n *= dimlen;
+ blk_size += sizeof (void *) * n;
+ }
+ size_t data_rows_num = n;
+ size_t data_rows_offset = blk_size;
+ blk_size += elsize * n * dimlen;
+
+ void *blk = (void *) malloc (blk_size);
+ memset (blk, 0, blk_size);
+ void **curr_dim = (void **) blk;
+ n = 1;
+
+ for (int d = 0; d < ndims - 1; d++)
+ {
+ uintptr_t next_dim = (uintptr_t) (curr_dim + n * dimlen);
+ size_t next_dimlen = dimlen * (d < ndims - 2 ? sizeof (void *) : elsize);
+
+ for (int b = 0; b < n; b++)
+ for (int i = 0; i < dimlen; i++)
+ if (d < ndims - 1)
+ curr_dim[b * dimlen + i]
+ = (void*) (next_dim + b * dimlen * next_dimlen + i * next_dimlen);
+
+ n *= dimlen;
+ curr_dim = (void**) next_dim;
+ }
+ assert (n == data_rows_num);
+ return blk;
+}
^ permalink raw reply [flat|nested] 6+ messages in thread
* (v4 update) Re: [PATCH, OpenACC, v3] Non-contiguous array support for OpenACC data clauses
2019-11-26 14:57 [PATCH, OpenACC, v3] Non-contiguous array support for OpenACC data clauses Chung-Lin Tang
@ 2020-04-07 13:50 ` Chung-Lin Tang
2023-03-10 13:24 ` Thomas Schwinge
` (2 subsequent siblings)
3 siblings, 0 replies; 6+ messages in thread
From: Chung-Lin Tang @ 2020-04-07 13:50 UTC (permalink / raw)
To: gcc-patches, Thomas Schwinge; +Cc: Jakub Jelinek
[-- Attachment #1: Type: text/plain, Size: 5911 bytes --]
On 2019/11/26 10:49 PM, Chung-Lin Tang wrote:
> Hi Thomas,
> this is a reorg of the last non-contiguous arrays patch. You'll notice that:
>
> (1) A large part of the code has been pulled into oacc-parallel.c, with most
> of the data structure declarations in oacc-int.h.
>
> (2) target.c only contains relatively little code from gomp_map_vars_internal
> that processes what GOACC_parallel_keyed/data_start gives it.
>
> (3) Instead of directly passed in the map pointer, the array descriptor
> pointers are now passed to GOACC_parallel_keyed/data_start using varargs.
> (I believe the adding of '...' to GOACC_data_start does not break any
> compatiblity)
>
> (4) Along the way, I've added a 'gomp_map_vars_openacc' for specializing our
> uses, which should shave off quite some code through inlining.
>
> The GOMP_MAP_NONCONTIG_ARRAY_P maps are still placed at the beginning of the
> recieved map sequence in this patch. It should still be relatively easy to
> use a GOACC_FLAG_* to do so if deemed better before committing.
>
> Thanks,
> Chung-Lin
Hi Thomas,
this is a rebased version, with some updates WRT the attach/detach changes and
some bug fixes, dubbed "v4". Plan to merge this version to the OG10 branch soon.
Thanks,
Chung-Lin
> PR other/76739
>
> gcc/c/
> * c-typeck.c (handle_omp_array_sections_1): Add 'bool &non_contiguous'
> parameter, adjust recursive call site, add cases for allowing
> pointer based multi-dimensional arrays for OpenACC.
> (handle_omp_array_sections): Adjust handle_omp_array_sections_1 call,
> handle non-contiguous case to create dynamic array map.
>
> gcc/cp/
> * semantics.c (handle_omp_array_sections_1): Add 'bool &non_contiguous'
> parameter, adjust recursive call site, add cases for allowing
> pointer based multi-dimensional arrays for OpenACC.
> (handle_omp_array_sections): Adjust handle_omp_array_sections_1 call,
> handle non-contiguous case to create dynamic array map.
>
> gcc/fortran/
> * f95-lang.c (DEF_FUNCTION_TYPE_VAR_5): New symbol.
> * types.def (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR): New type.
>
> gcc/
> * builtin-types.def (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR): New type.
> * omp-builtins.def (BUILT_IN_GOACC_DATA_START): Adjust function type
> to new BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR.
> * gimplify.c (gimplify_scan_omp_clauses): Skip gimplification of
> OMP_CLAUSE_SIZE of non-contiguous array maps (which is a TREE_LIST).
> * omp-expand.c (expand_omp_target): Add non-contiguous array descriptor
> pointers to variadic arguments.
> * omp-low.c (append_field_to_record_type): New function.
> (create_noncontig_array_descr_type): Likewise.
> (create_noncontig_array_descr_init_code): Likewise.
> (scan_sharing_clauses): For non-contiguous array map kinds, check for
> supported dimension structure, and install non-contiguous array
> variable into current omp_context.
> (reorder_noncontig_array_clauses): New function.
> (scan_omp_target): Call reorder_noncontig_array_clauses to place
> non-contiguous array map clauses at beginning of clause sequence.
> (lower_omp_target): Add handling for non-contiguous array map kinds,
> add all created non-contiguous array descriptors to
> gimple_omp_target_data_arg.
>
> gcc/testsuite/
> * c-c++-common/goacc/noncontig_array-1.c: New test.
>
> libgomp/
> * libgomp_g.h (GOACC_data_start): Add variadic '...' to declaration.
> * libgomp.h (gomp_map_vars_openacc): New function declaration.
> * oacc-int.h (struct goacc_ncarray_dim): New struct declaration.
> (struct goacc_ncarray_descr_type): Likewise.
> (struct goacc_ncarray): Likewise.
> (struct goacc_ncarray_info): Likewise.
> (goacc_noncontig_array_create_ptrblock): New function declaration.
> * oacc-parallel.c (goacc_noncontig_array_count_rows): New function.
> (goacc_noncontig_array_compute_sizes): Likewise.
> (goacc_noncontig_array_fill_rows_1): Likewise.
> (goacc_noncontig_array_fill_rows): Likewise.
> (goacc_process_noncontiguous_arrays): Likewise.
> (goacc_noncontig_array_create_ptrblock): Likewise.
> (GOACC_parallel_keyed): Use goacc_process_noncontiguous_arrays to
> handle non-contiguous array descriptors at end of varargs, adjust
> to use gomp_map_vars_openacc.
> (GOACC_data_start): Likewise. Adjust function type to accept varargs.
> * target.c (gomp_map_vars_internal): Add struct goacc_ncarray_info *
> nca_info parameter, add handling code for non-contiguous arrays.
> (gomp_map_vars_openacc): Add new function for specialization of
> gomp_map_vars_internal for OpenACC structured region usage.
>
> * testsuite/libgomp.oacc-c-c++-common/noncontig_array-1.c: New test.
> * testsuite/libgomp.oacc-c-c++-common/noncontig_array-2.c: New test.
> * testsuite/libgomp.oacc-c-c++-common/noncontig_array-3.c: New test.
> * testsuite/libgomp.oacc-c-c++-common/noncontig_array-4.c: New test.
> * testsuite/libgomp.oacc-c-c++-common/noncontig_array-utils.h: Support
> header for new tests.
>
[-- Attachment #2: openacc-noncontig-arrays-v4.patch --]
[-- Type: text/plain, Size: 59558 bytes --]
diff --git a/gcc/builtin-types.def b/gcc/builtin-types.def
index c7aa691b243..5e36a498f4a 100644
--- a/gcc/builtin-types.def
+++ b/gcc/builtin-types.def
@@ -830,6 +830,9 @@ DEF_FUNCTION_TYPE_VAR_5 (BT_FN_INT_STRING_SIZE_INT_SIZE_CONST_STRING_VAR,
DEF_FUNCTION_TYPE_VAR_5 (BT_FN_INT_INT_INT_INT_INT_INT_VAR,
BT_INT, BT_INT, BT_INT, BT_INT, BT_INT, BT_INT)
+DEF_FUNCTION_TYPE_VAR_5 (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR,
+ BT_VOID, BT_INT, BT_SIZE, BT_PTR, BT_PTR, BT_PTR)
+
DEF_FUNCTION_TYPE_VAR_6 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_VAR,
BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE,
BT_PTR, BT_PTR, BT_PTR)
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index 385bf3a1c7b..192a2cc5281 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -12921,12 +12921,14 @@ c_finish_omp_cancellation_point (location_t loc, tree clauses)
<= FIRST_NON_ONE we diagnose non-contiguous arrays if low bound isn't
0 or length isn't the array domain max + 1, for > FIRST_NON_ONE we
can if MAYBE_ZERO_LEN is false. MAYBE_ZERO_LEN will be true in the above
- case though, as some lengths could be zero. */
+ case though, as some lengths could be zero.
+ NON_CONTIGUOUS will be true if this is an OpenACC non-contiguous array
+ section. */
static tree
handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
bool &maybe_zero_len, unsigned int &first_non_one,
- enum c_omp_region_type ort)
+ bool &non_contiguous, enum c_omp_region_type ort)
{
tree ret, low_bound, length, type;
if (TREE_CODE (t) != TREE_LIST)
@@ -13019,7 +13021,8 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
}
ret = handle_omp_array_sections_1 (c, TREE_CHAIN (t), types,
- maybe_zero_len, first_non_one, ort);
+ maybe_zero_len, first_non_one,
+ non_contiguous, ort);
if (ret == error_mark_node || ret == NULL_TREE)
return ret;
@@ -13237,14 +13240,21 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
return error_mark_node;
}
/* If there is a pointer type anywhere but in the very first
- array-section-subscript, the array section can't be contiguous. */
+ array-section-subscript, the array section can't be contiguous.
+ Note that OpenACC does accept these kinds of non-contiguous pointer
+ based arrays. */
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND
&& TREE_CODE (TREE_CHAIN (t)) == TREE_LIST)
{
- 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 (ort == C_ORT_ACC)
+ non_contiguous = true;
+ 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;
+ }
}
}
else
@@ -13273,6 +13283,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
{
bool maybe_zero_len = false;
unsigned int first_non_one = 0;
+ bool non_contiguous = false;
auto_vec<tree, 10> types;
tree *tp = &OMP_CLAUSE_DECL (c);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND
@@ -13282,7 +13293,7 @@ handle_omp_array_sections (tree c, 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,
- ort);
+ non_contiguous, ort);
if (first == error_mark_node)
return true;
if (first == NULL_TREE)
@@ -13315,6 +13326,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
unsigned int num = types.length (), i;
tree t, side_effects = NULL_TREE, size = NULL_TREE;
tree condition = NULL_TREE;
+ tree ncarray_dims = NULL_TREE;
if (int_size_in_bytes (TREE_TYPE (first)) <= 0)
maybe_zero_len = true;
@@ -13338,6 +13350,13 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
length = fold_convert (sizetype, length);
if (low_bound == NULL_TREE)
low_bound = integer_zero_node;
+
+ if (non_contiguous)
+ {
+ ncarray_dims = tree_cons (low_bound, length, ncarray_dims);
+ continue;
+ }
+
if (!maybe_zero_len && i > first_non_one)
{
if (integer_nonzerop (low_bound))
@@ -13434,6 +13453,14 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
size = size_binop (MULT_EXPR, size, l);
}
}
+ if (non_contiguous)
+ {
+ int kind = OMP_CLAUSE_MAP_KIND (c);
+ OMP_CLAUSE_SET_MAP_KIND (c, kind | GOMP_MAP_NONCONTIG_ARRAY);
+ OMP_CLAUSE_DECL (c) = t;
+ OMP_CLAUSE_SIZE (c) = ncarray_dims;
+ return false;
+ }
if (side_effects)
size = build2 (COMPOUND_EXPR, sizetype, side_effects, size);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index c7a6064e9f3..5649b771564 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -4737,12 +4737,14 @@ omp_privatize_field (tree t, bool shared)
<= FIRST_NON_ONE we diagnose non-contiguous arrays if low bound isn't
0 or length isn't the array domain max + 1, for > FIRST_NON_ONE we
can if MAYBE_ZERO_LEN is false. MAYBE_ZERO_LEN will be true in the above
- case though, as some lengths could be zero. */
+ case though, as some lengths could be zero.
+ NON_CONTIGUOUS will be true if this is an OpenACC non-contiguous array
+ section. */
static tree
handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
bool &maybe_zero_len, unsigned int &first_non_one,
- enum c_omp_region_type ort)
+ bool &non_contiguous, enum c_omp_region_type ort)
{
tree ret, low_bound, length, type;
if (TREE_CODE (t) != TREE_LIST)
@@ -4828,7 +4830,8 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
&& TREE_CODE (TREE_CHAIN (t)) == FIELD_DECL)
TREE_CHAIN (t) = omp_privatize_field (TREE_CHAIN (t), false);
ret = handle_omp_array_sections_1 (c, TREE_CHAIN (t), types,
- maybe_zero_len, first_non_one, ort);
+ maybe_zero_len, first_non_one,
+ non_contiguous, ort);
if (ret == error_mark_node || ret == NULL_TREE)
return ret;
@@ -5059,14 +5062,21 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
return error_mark_node;
}
/* If there is a pointer type anywhere but in the very first
- array-section-subscript, the array section can't be contiguous. */
+ array-section-subscript, the array section can't be contiguous.
+ Note that OpenACC does accept these kinds of non-contiguous pointer
+ based arrays. */
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND
&& TREE_CODE (TREE_CHAIN (t)) == TREE_LIST)
{
- 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 (ort == C_ORT_ACC)
+ non_contiguous = true;
+ 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;
+ }
}
}
else
@@ -5106,6 +5116,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
{
bool maybe_zero_len = false;
unsigned int first_non_one = 0;
+ bool non_contiguous = false;
auto_vec<tree, 10> types;
tree *tp = &OMP_CLAUSE_DECL (c);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND
@@ -5115,7 +5126,7 @@ handle_omp_array_sections (tree c, 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,
- ort);
+ non_contiguous, ort);
if (first == error_mark_node)
return true;
if (first == NULL_TREE)
@@ -5149,6 +5160,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
unsigned int num = types.length (), i;
tree t, side_effects = NULL_TREE, size = NULL_TREE;
tree condition = NULL_TREE;
+ tree ncarray_dims = NULL_TREE;
if (int_size_in_bytes (TREE_TYPE (first)) <= 0)
maybe_zero_len = true;
@@ -5174,6 +5186,13 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
length = fold_convert (sizetype, length);
if (low_bound == NULL_TREE)
low_bound = integer_zero_node;
+
+ if (non_contiguous)
+ {
+ ncarray_dims = tree_cons (low_bound, length, ncarray_dims);
+ continue;
+ }
+
if (!maybe_zero_len && i > first_non_one)
{
if (integer_nonzerop (low_bound))
@@ -5265,6 +5284,14 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
}
if (!processing_template_decl)
{
+ if (non_contiguous)
+ {
+ int kind = OMP_CLAUSE_MAP_KIND (c);
+ OMP_CLAUSE_SET_MAP_KIND (c, kind | GOMP_MAP_NONCONTIG_ARRAY);
+ OMP_CLAUSE_DECL (c) = t;
+ OMP_CLAUSE_SIZE (c) = ncarray_dims;
+ return false;
+ }
if (side_effects)
size = build2 (COMPOUND_EXPR, sizetype, side_effects, size);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION
diff --git a/gcc/fortran/f95-lang.c b/gcc/fortran/f95-lang.c
index 44ebe3e294d..a13dd7ee1cf 100644
--- a/gcc/fortran/f95-lang.c
+++ b/gcc/fortran/f95-lang.c
@@ -632,6 +632,8 @@ gfc_init_builtin_functions (void)
#define DEF_FUNCTION_TYPE_VAR_0(NAME, RETURN) NAME,
#define DEF_FUNCTION_TYPE_VAR_1(NAME, RETURN, ARG1) NAME,
#define DEF_FUNCTION_TYPE_VAR_2(NAME, RETURN, ARG1, ARG2) NAME,
+#define DEF_FUNCTION_TYPE_VAR_5(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5) \
+ NAME,
#define DEF_FUNCTION_TYPE_VAR_6(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
ARG6) NAME,
#define DEF_FUNCTION_TYPE_VAR_7(NAME, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
@@ -1145,6 +1147,15 @@ gfc_init_builtin_functions (void)
builtin_types[(int) ARG1], \
builtin_types[(int) ARG2], \
NULL_TREE);
+#define DEF_FUNCTION_TYPE_VAR_5(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5) \
+ builtin_types[(int) ENUM] \
+ = build_varargs_function_type_list (builtin_types[(int) RETURN], \
+ builtin_types[(int) ARG1], \
+ builtin_types[(int) ARG2], \
+ builtin_types[(int) ARG3], \
+ builtin_types[(int) ARG4], \
+ builtin_types[(int) ARG5], \
+ NULL_TREE);
#define DEF_FUNCTION_TYPE_VAR_6(ENUM, RETURN, ARG1, ARG2, ARG3, ARG4, ARG5, \
ARG6) \
builtin_types[(int) ENUM] \
@@ -1186,6 +1197,7 @@ gfc_init_builtin_functions (void)
#undef DEF_FUNCTION_TYPE_VAR_0
#undef DEF_FUNCTION_TYPE_VAR_1
#undef DEF_FUNCTION_TYPE_VAR_2
+#undef DEF_FUNCTION_TYPE_VAR_5
#undef DEF_FUNCTION_TYPE_VAR_6
#undef DEF_FUNCTION_TYPE_VAR_7
#undef DEF_POINTER_TYPE
diff --git a/gcc/fortran/types.def b/gcc/fortran/types.def
index 7b4925cdd7e..e5fafc5a34d 100644
--- a/gcc/fortran/types.def
+++ b/gcc/fortran/types.def
@@ -269,6 +269,9 @@ DEF_FUNCTION_TYPE_VAR_7 (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_INT_INT_VAR,
BT_VOID, BT_INT, BT_SIZE, BT_PTR, BT_PTR,
BT_PTR, BT_INT, BT_INT)
+DEF_FUNCTION_TYPE_VAR_5 (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR,
+ BT_VOID, BT_INT, BT_SIZE, BT_PTR, BT_PTR, BT_PTR)
+
DEF_FUNCTION_TYPE_VAR_6 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_VAR,
BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE,
BT_PTR, BT_PTR, BT_PTR)
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 787435c38cd..0059167a441 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -8810,8 +8810,16 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
OMP_CLAUSE_SIZE (c) = DECL_P (decl) ? DECL_SIZE_UNIT (decl)
: TYPE_SIZE_UNIT (TREE_TYPE (decl));
- if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p,
- NULL, is_gimple_val, fb_rvalue) == GS_ERROR)
+ if (GOMP_MAP_NONCONTIG_ARRAY_P (OMP_CLAUSE_MAP_KIND (c)))
+ {
+ gcc_assert (OMP_CLAUSE_SIZE (c)
+ && TREE_CODE (OMP_CLAUSE_SIZE (c)) == TREE_LIST);
+ /* For non-contiguous array maps, OMP_CLAUSE_SIZE is a TREE_LIST
+ of the individual array dimensions, which gimplify_expr doesn't
+ handle, so skip the call to gimplify_expr here. */
+ }
+ else if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p,
+ NULL, is_gimple_val, fb_rvalue) == GS_ERROR)
{
remove = true;
break;
diff --git a/gcc/omp-builtins.def b/gcc/omp-builtins.def
index f461d60e52b..ed8584cdc1c 100644
--- a/gcc/omp-builtins.def
+++ b/gcc/omp-builtins.def
@@ -32,7 +32,7 @@ along with GCC; see the file COPYING3. If not see
DEF_GOACC_BUILTIN (BUILT_IN_ACC_GET_DEVICE_TYPE, "acc_get_device_type",
BT_FN_INT, ATTR_NOTHROW_LIST)
DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DATA_START, "GOACC_data_start",
- BT_FN_VOID_INT_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
+ BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR, ATTR_NOTHROW_LIST)
DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DATA_END, "GOACC_data_end",
BT_FN_VOID, ATTR_NOTHROW_LIST)
DEF_GOACC_BUILTIN (BUILT_IN_GOACC_ENTER_EXIT_DATA, "GOACC_enter_exit_data",
diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c
index a642ccc9980..2951b1ae1ec 100644
--- a/gcc/omp-expand.c
+++ b/gcc/omp-expand.c
@@ -8459,6 +8459,19 @@ expand_omp_target (struct omp_region *region)
/* Push terminal marker - zero. */
args.safe_push (oacc_launch_pack (0, NULL_TREE, 0));
+ /* We assume index >= 3 in gimple_omp_target_data_arg are non-contiguous
+ array descriptor pointer arguments. */
+ if (t != NULL
+ && TREE_VEC_LENGTH (t) > 3
+ && (start_ix == BUILT_IN_GOACC_DATA_START
+ || start_ix == BUILT_IN_GOACC_PARALLEL))
+ {
+ gcc_assert ((c = omp_find_clause (clauses, OMP_CLAUSE_MAP))
+ && GOMP_MAP_NONCONTIG_ARRAY_P (OMP_CLAUSE_MAP_KIND (c)));
+ for (int i = 3; i < TREE_VEC_LENGTH (t); i++)
+ args.safe_push (TREE_VEC_ELT (t, i));
+ }
+
g = gimple_build_call_vec (builtin_decl_explicit (start_ix), args);
gimple_set_location (g, gimple_location (entry_stmt));
gsi_insert_before (&gsi, g, GSI_SAME_STMT);
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 67565d61400..dcb34f99ef1 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -915,6 +915,123 @@ omp_copy_decl (tree var, copy_body_data *cb)
return error_mark_node;
}
+/* Helper function for create_noncontig_array_descr_type(), to append a new field
+ to a record type. */
+
+static void
+append_field_to_record_type (tree record_type, tree fld_ident, tree fld_type)
+{
+ tree *p, fld = build_decl (UNKNOWN_LOCATION, FIELD_DECL, fld_ident, fld_type);
+ DECL_CONTEXT (fld) = record_type;
+
+ for (p = &TYPE_FIELDS (record_type); *p; p = &DECL_CHAIN (*p))
+ ;
+ *p = fld;
+}
+
+/* Create type for non-contiguous array descriptor. Returns created type, and
+ returns the number of dimensions in *DIM_NUM. */
+
+static tree
+create_noncontig_array_descr_type (tree dims, int *dim_num)
+{
+ int n = 0;
+ tree array_descr_type, name, x;
+ gcc_assert (TREE_CODE (dims) == TREE_LIST);
+
+ array_descr_type = lang_hooks.types.make_type (RECORD_TYPE);
+ name = create_tmp_var_name (".omp_noncontig_array_descr_type");
+ name = build_decl (UNKNOWN_LOCATION, TYPE_DECL, name, array_descr_type);
+ DECL_ARTIFICIAL (name) = 1;
+ DECL_NAMELESS (name) = 1;
+ TYPE_NAME (array_descr_type) = name;
+ TYPE_ARTIFICIAL (array_descr_type) = 1;
+
+ /* Number of dimensions. */
+ append_field_to_record_type (array_descr_type, get_identifier ("__dim_num"),
+ sizetype);
+
+ for (x = dims; x; x = TREE_CHAIN (x), n++)
+ {
+ char *fldname;
+ /* One for the start index. */
+ ASM_FORMAT_PRIVATE_NAME (fldname, "__dim_base", n);
+ append_field_to_record_type (array_descr_type, get_identifier (fldname),
+ sizetype);
+ /* One for the length. */
+ ASM_FORMAT_PRIVATE_NAME (fldname, "__dim_length", n);
+ append_field_to_record_type (array_descr_type, get_identifier (fldname),
+ sizetype);
+ /* One for the element size. */
+ ASM_FORMAT_PRIVATE_NAME (fldname, "__dim_elem_size", n);
+ append_field_to_record_type (array_descr_type, get_identifier (fldname),
+ sizetype);
+ /* One for is_array flag. */
+ ASM_FORMAT_PRIVATE_NAME (fldname, "__dim_is_array", n);
+ append_field_to_record_type (array_descr_type, get_identifier (fldname),
+ sizetype);
+ }
+
+ layout_type (array_descr_type);
+ *dim_num = n;
+ return array_descr_type;
+}
+
+/* Generate code sequence for initializing non-contiguous array descriptor. */
+
+static void
+create_noncontig_array_descr_init_code (tree array_descr, tree array_var,
+ tree dimensions, int dim_num,
+ gimple_seq *ilist)
+{
+ tree fld, fldref;
+ tree array_descr_type = TREE_TYPE (array_descr);
+ tree dim_type = TREE_TYPE (array_var);
+
+ if (TREE_CODE (dim_type) == REFERENCE_TYPE)
+ dim_type = TREE_TYPE (dim_type);
+
+ fld = TYPE_FIELDS (array_descr_type);
+ fldref = omp_build_component_ref (array_descr, fld);
+ gimplify_assign (fldref, build_int_cst (sizetype, dim_num), ilist);
+
+ while (dimensions)
+ {
+ tree dim_base = fold_convert (sizetype, TREE_PURPOSE (dimensions));
+ tree dim_length = fold_convert (sizetype, TREE_VALUE (dimensions));
+ tree dim_elem_size = TYPE_SIZE_UNIT (TREE_TYPE (dim_type));
+ tree dim_is_array = (TREE_CODE (dim_type) == ARRAY_TYPE
+ ? integer_one_node : integer_zero_node);
+ /* Set base. */
+ fld = TREE_CHAIN (fld);
+ fldref = omp_build_component_ref (array_descr, fld);
+ dim_base = fold_build2 (MULT_EXPR, sizetype, dim_base, dim_elem_size);
+ gimplify_assign (fldref, dim_base, ilist);
+
+ /* Set length. */
+ fld = TREE_CHAIN (fld);
+ fldref = omp_build_component_ref (array_descr, fld);
+ dim_length = fold_build2 (MULT_EXPR, sizetype, dim_length, dim_elem_size);
+ gimplify_assign (fldref, dim_length, ilist);
+
+ /* Set elem_size. */
+ fld = TREE_CHAIN (fld);
+ fldref = omp_build_component_ref (array_descr, fld);
+ dim_elem_size = fold_convert (sizetype, dim_elem_size);
+ gimplify_assign (fldref, dim_elem_size, ilist);
+
+ /* Set is_array flag. */
+ fld = TREE_CHAIN (fld);
+ fldref = omp_build_component_ref (array_descr, fld);
+ dim_is_array = fold_convert (sizetype, dim_is_array);
+ gimplify_assign (fldref, dim_is_array, ilist);
+
+ dimensions = TREE_CHAIN (dimensions);
+ dim_type = TREE_TYPE (dim_type);
+ }
+ gcc_assert (TREE_CHAIN (fld) == NULL_TREE);
+}
+
/* Create a new context, with OUTER_CTX being the surrounding context. */
static omp_context *
@@ -1397,6 +1514,38 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
install_var_local (decl, ctx);
break;
}
+
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && GOMP_MAP_NONCONTIG_ARRAY_P (OMP_CLAUSE_MAP_KIND (c)))
+ {
+ tree array_decl = OMP_CLAUSE_DECL (c);
+ tree array_type = TREE_TYPE (array_decl);
+ bool by_ref = (TREE_CODE (array_type) == ARRAY_TYPE
+ ? true : false);
+
+ /* Checking code to ensure we only have arrays at top dimension.
+ This limitation might be lifted in the future. See PR76639. */
+ if (TREE_CODE (array_type) == REFERENCE_TYPE)
+ array_type = TREE_TYPE (array_type);
+ tree t = array_type, prev_t = NULL_TREE;
+ while (t)
+ {
+ if (TREE_CODE (t) == ARRAY_TYPE && prev_t)
+ {
+ error_at (gimple_location (ctx->stmt), "array types are"
+ " only allowed at outermost dimension of"
+ " non-contiguous array");
+ break;
+ }
+ prev_t = t;
+ t = TREE_TYPE (t);
+ }
+
+ install_var_field (array_decl, by_ref, 3, ctx);
+ install_var_local (array_decl, ctx);
+ break;
+ }
+
if (DECL_P (decl))
{
if (DECL_SIZE (decl)
@@ -2709,6 +2858,50 @@ scan_omp_single (gomp_single *stmt, omp_context *outer_ctx)
layout_type (ctx->record_type);
}
+/* Reorder clauses so that non-contiguous array map clauses are placed at the very
+ front of the chain. */
+
+static void
+reorder_noncontig_array_clauses (tree *clauses_ptr)
+{
+ tree c, clauses = *clauses_ptr;
+ tree prev_clause = NULL_TREE, next_clause;
+ tree array_clauses = NULL_TREE, array_clauses_tail = NULL_TREE;
+
+ for (c = clauses; c; c = next_clause)
+ {
+ next_clause = OMP_CLAUSE_CHAIN (c);
+
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && GOMP_MAP_NONCONTIG_ARRAY_P (OMP_CLAUSE_MAP_KIND (c)))
+ {
+ /* Unchain c from clauses. */
+ if (c == clauses)
+ clauses = next_clause;
+
+ /* Link on to array_clauses. */
+ if (array_clauses_tail)
+ OMP_CLAUSE_CHAIN (array_clauses_tail) = c;
+ else
+ array_clauses = c;
+ array_clauses_tail = c;
+
+ if (prev_clause)
+ OMP_CLAUSE_CHAIN (prev_clause) = next_clause;
+ continue;
+ }
+
+ prev_clause = c;
+ }
+
+ /* Place non-contiguous array clauses at the start of the clause list. */
+ if (array_clauses)
+ {
+ OMP_CLAUSE_CHAIN (array_clauses_tail) = clauses;
+ *clauses_ptr = array_clauses;
+ }
+}
+
/* Scan a GIMPLE_OMP_TARGET. */
static void
@@ -2717,7 +2910,6 @@ scan_omp_target (gomp_target *stmt, omp_context *outer_ctx)
omp_context *ctx;
tree name;
bool offloaded = is_gimple_omp_offloaded (stmt);
- tree clauses = gimple_omp_target_clauses (stmt);
ctx = new_omp_context (stmt, outer_ctx);
ctx->field_map = splay_tree_new (splay_tree_compare_pointers, 0, 0);
@@ -2736,6 +2928,14 @@ scan_omp_target (gomp_target *stmt, omp_context *outer_ctx)
gimple_omp_target_set_child_fn (stmt, ctx->cb.dst_fn);
}
+ /* If is OpenACC construct, put non-contiguous array clauses (if any)
+ in front of clause chain. The runtime can then test the first to see
+ if the additional map processing for them is required. */
+ if (is_gimple_omp_oacc (stmt))
+ reorder_noncontig_array_clauses (gimple_omp_target_clauses_ptr (stmt));
+
+ tree clauses = gimple_omp_target_clauses (stmt);
+
scan_sharing_clauses (clauses, ctx);
scan_omp (gimple_omp_body_ptr (stmt), ctx);
@@ -11458,6 +11658,15 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
case GOMP_MAP_FORCE_PRESENT:
case GOMP_MAP_FORCE_DEVICEPTR:
case GOMP_MAP_DEVICE_RESIDENT:
+ case GOMP_MAP_NONCONTIG_ARRAY_TO:
+ case GOMP_MAP_NONCONTIG_ARRAY_FROM:
+ case GOMP_MAP_NONCONTIG_ARRAY_TOFROM:
+ case GOMP_MAP_NONCONTIG_ARRAY_FORCE_TO:
+ case GOMP_MAP_NONCONTIG_ARRAY_FORCE_FROM:
+ case GOMP_MAP_NONCONTIG_ARRAY_FORCE_TOFROM:
+ case GOMP_MAP_NONCONTIG_ARRAY_ALLOC:
+ case GOMP_MAP_NONCONTIG_ARRAY_FORCE_ALLOC:
+ case GOMP_MAP_NONCONTIG_ARRAY_FORCE_PRESENT:
case GOMP_MAP_LINK:
case GOMP_MAP_ATTACH:
case GOMP_MAP_DETACH:
@@ -11523,7 +11732,14 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
if (offloaded && !(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_IN_REDUCTION (c)))
{
- x = build_receiver_ref (var, true, ctx);
+ tree var_type = TREE_TYPE (var);
+ bool rcv_by_ref =
+ (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && GOMP_MAP_NONCONTIG_ARRAY_P (OMP_CLAUSE_MAP_KIND (c))
+ && TREE_CODE (var_type) != ARRAY_TYPE
+ ? false : true);
+
+ x = build_receiver_ref (var, rcv_by_ref, ctx);
tree new_var = lookup_decl (var, ctx);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
@@ -11697,6 +11913,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
vec_alloc (vkind, map_cnt);
unsigned int map_idx = 0;
+ vec<tree> nca_descrs = vNULL;
+
for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
switch (OMP_CLAUSE_CODE (c))
{
@@ -11773,6 +11991,29 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
avar = build_fold_addr_expr (avar);
gimplify_assign (x, avar, &ilist);
}
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && GOMP_MAP_NONCONTIG_ARRAY_P (OMP_CLAUSE_MAP_KIND (c)))
+ {
+ int dim_num;
+ tree dimensions = OMP_CLAUSE_SIZE (c);
+
+ tree array_descr_type =
+ create_noncontig_array_descr_type (dimensions, &dim_num);
+ tree array_descr =
+ create_tmp_var_raw (array_descr_type,
+ ".omp_noncontig_array_descr");
+ TREE_ADDRESSABLE (array_descr) = 1;
+ TREE_STATIC (array_descr) = 1;
+ gimple_add_tmp_var (array_descr);
+
+ create_noncontig_array_descr_init_code
+ (array_descr, ovar, dimensions, dim_num, &ilist);
+ nca_descrs.safe_push (build_fold_addr_expr (array_descr));
+
+ gimplify_assign (x, (TREE_CODE (TREE_TYPE (ovar)) == ARRAY_TYPE
+ ? build_fold_addr_expr (ovar) : ovar),
+ &ilist);
+ }
else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
{
gcc_assert (is_gimple_omp_oacc (ctx->stmt));
@@ -11845,6 +12086,9 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
s = TREE_TYPE (s);
s = TYPE_SIZE_UNIT (s);
}
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && GOMP_MAP_NONCONTIG_ARRAY_P (OMP_CLAUSE_MAP_KIND (c)))
+ s = NULL_TREE;
else
s = OMP_CLAUSE_SIZE (c);
if (s == NULL_TREE)
@@ -12097,6 +12341,19 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
gcc_assert (map_idx == map_cnt);
+ unsigned nca_num = nca_descrs.length ();
+ if (nca_num > 0)
+ {
+ tree nca, t = gimple_omp_target_data_arg (stmt);
+ int i, oldlen = TREE_VEC_LENGTH (t);
+ tree nt = make_tree_vec (oldlen + nca_num);
+ for (i = 0; i < oldlen; i++)
+ TREE_VEC_ELT (nt, i) = TREE_VEC_ELT (t, i);
+ for (i = 0; nca_descrs.iterate (i, &nca); i++)
+ TREE_VEC_ELT (nt, oldlen + i) = nca;
+ gimple_omp_target_set_data_arg (stmt, nt);
+ }
+
DECL_INITIAL (TREE_VEC_ELT (t, 1))
= build_constructor (TREE_TYPE (TREE_VEC_ELT (t, 1)), vsize);
DECL_INITIAL (TREE_VEC_ELT (t, 2))
diff --git a/gcc/testsuite/c-c++-common/goacc/noncontig_array-1.c b/gcc/testsuite/c-c++-common/goacc/noncontig_array-1.c
new file mode 100644
index 00000000000..ea738f5b65b
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/noncontig_array-1.c
@@ -0,0 +1,25 @@
+/* { dg-do compile } */
+
+void foo (void)
+{
+ int array_of_array[10][10];
+ int **ptr_to_ptr;
+ int *array_of_ptr[10];
+ int (*ptr_to_array)[10];
+
+ #pragma acc parallel copy (array_of_array[2:4][0:10])
+ array_of_array[5][5] = 1;
+
+ #pragma acc parallel copy (ptr_to_ptr[2:4][1:7])
+ ptr_to_ptr[5][5] = 1;
+
+ #pragma acc parallel copy (array_of_ptr[2:4][1:7])
+ array_of_ptr[5][5] = 1;
+
+ #pragma acc parallel copy (ptr_to_array[2:4][1:7]) /* { dg-error "array section is not contiguous in 'map' clause" } */
+ ptr_to_array[5][5] = 1;
+}
+/* { dg-final { scan-tree-dump-times {#pragma omp target oacc_parallel map\(tofrom:array_of_array} 1 gimple } } */
+/* { dg-final { scan-tree-dump-times {#pragma omp target oacc_parallel map\(tofrom,noncontig_array:ptr_to_ptr \[dimensions: 2 4, 1 7\]} 1 gimple } } */
+/* { dg-final { scan-tree-dump-times {#pragma omp target oacc_parallel map\(tofrom,noncontig_array:array_of_ptr \[dimensions: 2 4, 1 7\]} 1 gimple } } */
+/* { dg-final { scan-tree-dump-times {#pragma omp target oacc_parallel map\(tofrom,noncontig_array:ptr_to_array \[dimensions: 2 4, 1 7\]} 1 gimple { xfail *-*-* } } } */
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index 885ca8cd329..6ceafa399df 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -863,6 +863,33 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
case GOMP_MAP_ATTACH_DETACH:
pp_string (pp, "attach_detach");
break;
+ case GOMP_MAP_NONCONTIG_ARRAY_TO:
+ pp_string (pp, "to,noncontig_array");
+ break;
+ case GOMP_MAP_NONCONTIG_ARRAY_FROM:
+ pp_string (pp, "from,noncontig_array");
+ break;
+ case GOMP_MAP_NONCONTIG_ARRAY_TOFROM:
+ pp_string (pp, "tofrom,noncontig_array");
+ break;
+ case GOMP_MAP_NONCONTIG_ARRAY_FORCE_TO:
+ pp_string (pp, "force_to,noncontig_array");
+ break;
+ case GOMP_MAP_NONCONTIG_ARRAY_FORCE_FROM:
+ pp_string (pp, "force_from,noncontig_array");
+ break;
+ case GOMP_MAP_NONCONTIG_ARRAY_FORCE_TOFROM:
+ pp_string (pp, "force_tofrom,noncontig_array");
+ break;
+ case GOMP_MAP_NONCONTIG_ARRAY_ALLOC:
+ pp_string (pp, "alloc,noncontig_array");
+ break;
+ case GOMP_MAP_NONCONTIG_ARRAY_FORCE_ALLOC:
+ pp_string (pp, "force_alloc,noncontig_array");
+ break;
+ case GOMP_MAP_NONCONTIG_ARRAY_FORCE_PRESENT:
+ pp_string (pp, "force_present,noncontig_array");
+ break;
default:
gcc_unreachable ();
}
@@ -873,8 +900,15 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
if (OMP_CLAUSE_SIZE (clause))
{
switch (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP
- ? OMP_CLAUSE_MAP_KIND (clause) : GOMP_MAP_TO)
+ ? (GOMP_MAP_NONCONTIG_ARRAY_P (OMP_CLAUSE_MAP_KIND (clause))
+ ? GOMP_MAP_NONCONTIG_ARRAY
+ : OMP_CLAUSE_MAP_KIND (clause))
+ : GOMP_MAP_TO)
{
+ case GOMP_MAP_NONCONTIG_ARRAY:
+ gcc_assert (TREE_CODE (OMP_CLAUSE_SIZE (clause)) == TREE_LIST);
+ pp_string (pp, " [dimensions: ");
+ break;
case GOMP_MAP_POINTER:
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
diff --git a/include/gomp-constants.h b/include/gomp-constants.h
index 1587e4d2ba2..927b65c2c7d 100644
--- a/include/gomp-constants.h
+++ b/include/gomp-constants.h
@@ -40,6 +40,7 @@
#define GOMP_MAP_FLAG_SPECIAL_0 (1 << 2)
#define GOMP_MAP_FLAG_SPECIAL_1 (1 << 3)
#define GOMP_MAP_FLAG_SPECIAL_2 (1 << 4)
+#define GOMP_MAP_FLAG_SPECIAL_3 (1 << 5)
#define GOMP_MAP_FLAG_SPECIAL_4 (1 << 6)
#define GOMP_MAP_FLAG_SPECIAL (GOMP_MAP_FLAG_SPECIAL_1 \
| GOMP_MAP_FLAG_SPECIAL_0)
@@ -143,6 +144,26 @@ enum gomp_map_kind
/* In OpenACC, detach a pointer to a mapped struct field. */
GOMP_MAP_FORCE_DETACH = (GOMP_MAP_DEEP_COPY
| GOMP_MAP_FLAG_FORCE | 1),
+ /* Mapping kinds for non-contiguous arrays. */
+ GOMP_MAP_NONCONTIG_ARRAY = (GOMP_MAP_FLAG_SPECIAL_3),
+ GOMP_MAP_NONCONTIG_ARRAY_TO = (GOMP_MAP_NONCONTIG_ARRAY
+ | GOMP_MAP_TO),
+ GOMP_MAP_NONCONTIG_ARRAY_FROM = (GOMP_MAP_NONCONTIG_ARRAY
+ | GOMP_MAP_FROM),
+ GOMP_MAP_NONCONTIG_ARRAY_TOFROM = (GOMP_MAP_NONCONTIG_ARRAY
+ | GOMP_MAP_TOFROM),
+ GOMP_MAP_NONCONTIG_ARRAY_FORCE_TO = (GOMP_MAP_NONCONTIG_ARRAY_TO
+ | GOMP_MAP_FLAG_FORCE),
+ GOMP_MAP_NONCONTIG_ARRAY_FORCE_FROM = (GOMP_MAP_NONCONTIG_ARRAY_FROM
+ | GOMP_MAP_FLAG_FORCE),
+ GOMP_MAP_NONCONTIG_ARRAY_FORCE_TOFROM = (GOMP_MAP_NONCONTIG_ARRAY_TOFROM
+ | GOMP_MAP_FLAG_FORCE),
+ GOMP_MAP_NONCONTIG_ARRAY_ALLOC = (GOMP_MAP_NONCONTIG_ARRAY
+ | GOMP_MAP_ALLOC),
+ GOMP_MAP_NONCONTIG_ARRAY_FORCE_ALLOC = (GOMP_MAP_NONCONTIG_ARRAY
+ | GOMP_MAP_FORCE_ALLOC),
+ GOMP_MAP_NONCONTIG_ARRAY_FORCE_PRESENT = (GOMP_MAP_NONCONTIG_ARRAY
+ | GOMP_MAP_FORCE_PRESENT),
/* Internal to GCC, not used in libgomp. */
/* Do not map, but pointer assign a pointer instead. */
@@ -175,6 +196,8 @@ enum gomp_map_kind
#define GOMP_MAP_ALWAYS_P(X) \
(GOMP_MAP_ALWAYS_TO_P (X) || ((X) == GOMP_MAP_ALWAYS_FROM))
+#define GOMP_MAP_NONCONTIG_ARRAY_P(X) \
+ ((X) & GOMP_MAP_NONCONTIG_ARRAY)
/* Asynchronous behavior. Keep in sync with
libgomp/{openacc.h,openacc.f90,openacc_lib.h}:acc_async_t. */
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index f5415bb156c..2acf3e3f920 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -1182,6 +1182,10 @@ extern struct target_mem_desc *gomp_map_vars_async (struct gomp_device_descr *,
size_t, void **, void **,
size_t *, void *, bool,
enum gomp_map_vars_kind);
+extern struct target_mem_desc *gomp_map_vars_openacc (struct gomp_device_descr *,
+ struct goacc_asyncqueue *,
+ size_t, void **, size_t *,
+ unsigned short *, void *);
extern void gomp_unmap_vars (struct target_mem_desc *, bool);
extern void gomp_unmap_vars_async (struct target_mem_desc *, bool,
struct goacc_asyncqueue *);
diff --git a/libgomp/libgomp_g.h b/libgomp/libgomp_g.h
index 59e3697bfd8..70a27572182 100644
--- a/libgomp/libgomp_g.h
+++ b/libgomp/libgomp_g.h
@@ -373,7 +373,7 @@ extern void GOACC_parallel_keyed (int, void (*) (void *), size_t,
extern void GOACC_parallel (int, void (*) (void *), size_t, void **, size_t *,
unsigned short *, int, int, int, int, int, ...);
extern void GOACC_data_start (int, size_t, void **, size_t *,
- unsigned short *);
+ unsigned short *, ...);
extern void GOACC_data_end (void);
extern void GOACC_update (int, size_t, void **, size_t *,
unsigned short *, int, int, ...);
diff --git a/libgomp/oacc-int.h b/libgomp/oacc-int.h
index 3c2c9b84b2f..6a4879de909 100644
--- a/libgomp/oacc-int.h
+++ b/libgomp/oacc-int.h
@@ -165,6 +165,57 @@ bool _goacc_profiling_setup_p (struct goacc_thread *,
void goacc_profiling_dispatch (acc_prof_info *, acc_event_info *,
acc_api_info *);
+/* Definitions for data structures describing OpenACC non-contiguous arrays
+ (Note: interfaces with compiler)
+
+ The compiler generates a descriptor for each such array, places the
+ descriptor on stack, and passes the address of the descriptor to the libgomp
+ runtime as a normal map argument. The runtime then processes the array
+ data structure setup, and replaces the argument with the new actual
+ array address for the child function.
+
+ Care must be taken such that the struct field and layout assumptions
+ of struct goacc_ncarray_dim, goacc_ncarray_descr_type inside the compiler
+ be consistant with the below declarations. */
+
+struct goacc_ncarray_dim {
+ size_t base;
+ size_t length;
+ size_t elem_size;
+ size_t is_array;
+};
+
+struct goacc_ncarray_descr_type
+{
+ size_t ndims;
+ struct goacc_ncarray_dim dims[];
+};
+
+/* Internal non-contiguous array info struct, used only here inside the runtime. */
+
+struct goacc_ncarray
+{
+ struct goacc_ncarray_descr_type *descr;
+ void *ptr;
+ size_t map_index;
+ size_t ptrblock_size;
+ void **data_rows;
+ void **tgt_data_rows;
+ size_t data_row_num;
+ size_t data_row_size;
+};
+
+struct goacc_ncarray_info
+{
+ size_t num_data_rows, num_ncarray;
+ void **data_rows;
+ void **tgt_data_rows;
+ struct goacc_ncarray ncarray[];
+};
+
+extern void *goacc_noncontig_array_create_ptrblock (struct goacc_ncarray *, void *);
+
+
#ifdef HAVE_ATTRIBUTE_VISIBILITY
# pragma GCC visibility pop
#endif
diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c
index c7e46e35bd6..ba208fabdf7 100644
--- a/libgomp/oacc-parallel.c
+++ b/libgomp/oacc-parallel.c
@@ -36,7 +36,7 @@
#include <string.h>
#include <stdarg.h>
#include <assert.h>
-
+#include <stdio.h>
/* In the ABI, the GOACC_FLAGs are encoded as an inverted bitmask, so that we
continue to support the following two legacy values. */
@@ -92,6 +92,172 @@ handle_ftn_pointers (size_t mapnum, void **hostaddrs, size_t *sizes,
}
}
+static size_t
+goacc_noncontig_array_count_rows (struct goacc_ncarray_descr_type *descr)
+{
+ size_t nrows = 1;
+ for (size_t d = 0; d < descr->ndims - 1; d++)
+ nrows *= descr->dims[d].length / sizeof (void *);
+ return nrows;
+}
+
+static void
+goacc_noncontig_array_compute_sizes (struct goacc_ncarray *nca)
+{
+ size_t d, n = 1;
+ struct goacc_ncarray_descr_type *descr = nca->descr;
+
+ nca->ptrblock_size = 0;
+ for (d = 0; d < descr->ndims - 1; d++)
+ {
+ size_t dim_count = descr->dims[d].length / descr->dims[d].elem_size;
+ size_t dim_ptrblock_size = (descr->dims[d + 1].is_array
+ ? 0 : descr->dims[d].length * n);
+ nca->ptrblock_size += dim_ptrblock_size;
+ n *= dim_count;
+ }
+ nca->data_row_num = n;
+ nca->data_row_size = descr->dims[d].length;
+}
+
+static void
+goacc_noncontig_array_fill_rows_1 (struct goacc_ncarray_descr_type *descr, void *nca,
+ size_t d, void ***row_ptr, size_t *count)
+{
+ if (d < descr->ndims - 1)
+ {
+ size_t elsize = descr->dims[d].elem_size;
+ size_t n = descr->dims[d].length / elsize;
+ void *p = nca + descr->dims[d].base;
+ for (size_t i = 0; i < n; i++)
+ {
+ void *ptr = p + i * elsize;
+ /* Deref if next dimension is not array. */
+ if (!descr->dims[d + 1].is_array)
+ ptr = *((void **) ptr);
+ goacc_noncontig_array_fill_rows_1 (descr, ptr, d + 1, row_ptr, count);
+ }
+ }
+ else
+ {
+ **row_ptr = nca + descr->dims[d].base;
+ *row_ptr += 1;
+ *count += 1;
+ }
+}
+
+static size_t
+goacc_noncontig_array_fill_rows (struct goacc_ncarray *nca)
+{
+ size_t count = 0;
+ void **p = nca->data_rows;
+ goacc_noncontig_array_fill_rows_1 (nca->descr, nca->ptr, 0, &p, &count);
+ return count;
+}
+
+static struct goacc_ncarray_info *
+goacc_process_noncontiguous_arrays (size_t mapnum, void **hostaddrs,
+ unsigned short *kinds, va_list* ap)
+{
+ size_t i, nr, num_data_rows = 0, num_ncarray = 0, curr_row_start = 0;
+ struct goacc_ncarray_descr_type *descr;
+
+ /* We need to go over *ap twice, so preserve *ap state here. */
+ va_list itr;
+ va_copy (itr, *ap);
+ for (i = 0; i < mapnum; i++)
+ if (GOMP_MAP_NONCONTIG_ARRAY_P (kinds[i] & 0xff))
+ {
+ descr = va_arg (itr, struct goacc_ncarray_descr_type *);
+ num_data_rows += goacc_noncontig_array_count_rows (descr);
+ num_ncarray += 1;
+ }
+ else
+ break;
+
+ /* Allocate the entire info struct, array entries, and row pointer
+ arrays in one large block. */
+ struct goacc_ncarray_info *nca_info
+ = gomp_malloc (sizeof (struct goacc_ncarray_info)
+ + sizeof (struct goacc_ncarray) * num_ncarray
+ + sizeof (void *) * num_data_rows * 2);
+ nca_info->num_data_rows = num_data_rows;
+ nca_info->num_ncarray = num_ncarray;
+ nca_info->data_rows = (void **) (nca_info->ncarray + num_ncarray);
+ nca_info->tgt_data_rows = nca_info->data_rows + num_data_rows;
+
+ struct goacc_ncarray *curr_ncarray = nca_info->ncarray;
+ for (i = 0; i < mapnum; i++)
+ if (GOMP_MAP_NONCONTIG_ARRAY_P (kinds[i] & 0xff))
+ {
+ descr = va_arg (*ap, struct goacc_ncarray_descr_type *);
+ curr_ncarray->descr = descr;
+ curr_ncarray->ptr = hostaddrs[i];
+ curr_ncarray->map_index = i;
+
+ goacc_noncontig_array_compute_sizes (curr_ncarray);
+
+ curr_ncarray->data_rows = nca_info->data_rows + curr_row_start;
+ curr_ncarray->tgt_data_rows = nca_info->tgt_data_rows + curr_row_start;
+
+ nr = goacc_noncontig_array_fill_rows (curr_ncarray);
+ assert (nr == curr_ncarray->data_row_num);
+ curr_row_start += nr;
+ curr_ncarray += 1;
+ }
+ else
+ break;
+
+ return nca_info;
+}
+
+void *
+goacc_noncontig_array_create_ptrblock (struct goacc_ncarray *nca,
+ void *tgt_ptrblock_addr)
+{
+ struct goacc_ncarray_descr_type *descr = nca->descr;
+ void **tgt_data_rows = nca->tgt_data_rows;
+ void *ptrblock = gomp_malloc (nca->ptrblock_size);
+ void **curr_dim_ptrblock = (void **) ptrblock;
+ size_t n = 1;
+
+ for (size_t d = 0; d < descr->ndims - 1; d++)
+ {
+ int curr_dim_len = descr->dims[d].length;
+ int next_dim_len = descr->dims[d + 1].length;
+ int curr_dim_num = curr_dim_len / sizeof (void *);
+ size_t next_dim_bias = descr->dims[d + 1].base;
+
+ void *next_dim_ptrblock
+ = (void *)(curr_dim_ptrblock + n * curr_dim_num);
+
+ for (int b = 0; b < n; b++)
+ for (int i = 0; i < curr_dim_num; i++)
+ {
+ if (d < descr->ndims - 2)
+ {
+ void *ptr = (next_dim_ptrblock
+ + b * curr_dim_num * next_dim_len
+ + i * next_dim_len);
+ void *tgt_ptr = (tgt_ptrblock_addr
+ + (ptr - ptrblock) - next_dim_bias);
+ curr_dim_ptrblock[b * curr_dim_num + i] = tgt_ptr;
+ }
+ else
+ {
+ curr_dim_ptrblock[b * curr_dim_num + i]
+ = tgt_data_rows[b * curr_dim_num + i] - next_dim_bias;
+ }
+ void *addr = &curr_dim_ptrblock[b * curr_dim_num + i];
+ assert (ptrblock <= addr && addr < ptrblock + nca->ptrblock_size);
+ }
+
+ n *= curr_dim_num;
+ curr_dim_ptrblock = next_dim_ptrblock;
+ }
+ assert (n == nca->data_row_num);
+ return ptrblock;
+}
/* Launch a possibly offloaded function with FLAGS. FN is the host fn
address. MAPNUM, HOSTADDRS, SIZES & KINDS describe the memory
@@ -117,6 +283,7 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (void *),
int async = GOMP_ASYNC_SYNC;
unsigned dims[GOMP_DIM_MAX];
unsigned tag;
+ struct goacc_ncarray_info *nca_info = NULL;
#ifdef HAVE_INTTYPES_H
gomp_debug (0, "%s: mapnum=%"PRIu64", hostaddrs=%p, size=%p, kinds=%p\n",
@@ -249,11 +416,20 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (void *),
break;
}
+ /*case GOMP_LAUNCH_NONCONTIG_ARRAYS:
+ nca_info = goacc_process_noncontiguous_arrays (mapnum, hostaddrs,
+ kinds, &ap);
+ break;*/
+
default:
gomp_fatal ("unrecognized offload code '%d',"
" libgomp is too old", GOMP_LAUNCH_CODE (tag));
}
}
+
+ if (mapnum > 0 && GOMP_MAP_NONCONTIG_ARRAY_P (kinds[0] & 0xff))
+ nca_info = goacc_process_noncontiguous_arrays (mapnum, hostaddrs, kinds, &ap);
+
va_end (ap);
if (!(acc_dev->capabilities & GOMP_OFFLOAD_CAP_NATIVE_EXEC))
@@ -290,8 +466,10 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (void *),
goacc_aq aq = get_goacc_asyncqueue (async);
- tgt = gomp_map_vars_async (acc_dev, aq, mapnum, hostaddrs, NULL, sizes, kinds,
- true, GOMP_MAP_VARS_OPENACC);
+ tgt = gomp_map_vars_openacc (acc_dev, aq, mapnum, hostaddrs, sizes, kinds,
+ nca_info);
+ free (nca_info);
+
if (profiling_p)
{
prof_info.event_type = acc_ev_enter_data_end;
@@ -365,7 +543,7 @@ GOACC_parallel (int flags_m, void (*fn) (void *),
void
GOACC_data_start (int flags_m, size_t mapnum,
- void **hostaddrs, size_t *sizes, unsigned short *kinds)
+ void **hostaddrs, size_t *sizes, unsigned short *kinds, ...)
{
int flags = GOACC_FLAGS_UNMARSHAL (flags_m);
@@ -456,17 +634,26 @@ GOACC_data_start (int flags_m, size_t mapnum,
{
prof_info.device_type = acc_device_host;
api_info.device_type = prof_info.device_type;
- tgt = gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, true,
- GOMP_MAP_VARS_OPENACC);
+ tgt = gomp_map_vars_openacc (NULL, NULL, 0, NULL, NULL, NULL, NULL);
tgt->prev = thr->mapped_data;
thr->mapped_data = tgt;
goto out_prof;
}
+ struct goacc_ncarray_info *nca_info = NULL;
+ if (mapnum > 0 && GOMP_MAP_NONCONTIG_ARRAY_P (kinds[0] & 0xff))
+ {
+ va_list ap;
+ va_start (ap, kinds);
+ nca_info = goacc_process_noncontiguous_arrays (mapnum, hostaddrs, kinds, &ap);
+ va_end (ap);
+ }
+
gomp_debug (0, " %s: prepare mappings\n", __FUNCTION__);
- tgt = gomp_map_vars (acc_dev, mapnum, hostaddrs, NULL, sizes, kinds, true,
- GOMP_MAP_VARS_OPENACC);
+ tgt = gomp_map_vars_openacc (acc_dev, NULL, mapnum, hostaddrs, sizes, kinds,
+ nca_info);
+ free (nca_info);
gomp_debug (0, " %s: mappings prepared\n", __FUNCTION__);
tgt->prev = thr->mapped_data;
thr->mapped_data = tgt;
diff --git a/libgomp/target.c b/libgomp/target.c
index 36425477dcb..ca91fd85b20 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -654,18 +654,20 @@ static inline __attribute__((always_inline)) struct target_mem_desc *
gomp_map_vars_internal (struct gomp_device_descr *devicep,
struct goacc_asyncqueue *aq, size_t mapnum,
void **hostaddrs, void **devaddrs, size_t *sizes,
- void *kinds, bool short_mapkind,
- enum gomp_map_vars_kind pragma_kind)
+ void *kinds, struct goacc_ncarray_info *nca_info,
+ bool short_mapkind, enum gomp_map_vars_kind pragma_kind)
{
size_t i, tgt_align, tgt_size, not_found_cnt = 0;
+ size_t nca_data_row_num = (nca_info ? nca_info->num_data_rows : 0);
bool has_firstprivate = false;
const int rshift = short_mapkind ? 8 : 3;
const int typemask = short_mapkind ? 0xff : 0x7;
struct splay_tree_s *mem_map = &devicep->mem_map;
struct splay_tree_key_s cur_node;
struct target_mem_desc *tgt
- = gomp_malloc (sizeof (*tgt) + sizeof (tgt->list[0]) * mapnum);
- tgt->list_count = mapnum;
+ = gomp_malloc (sizeof (*tgt)
+ + sizeof (tgt->list[0]) * (mapnum + nca_data_row_num));
+ tgt->list_count = mapnum + nca_data_row_num;
tgt->refcount = (pragma_kind == GOMP_MAP_VARS_ENTER_DATA
|| pragma_kind == GOMP_MAP_VARS_OPENACC_ENTER_DATA) ? 0 : 1;
tgt->device_descr = devicep;
@@ -814,6 +816,28 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
has_firstprivate = true;
continue;
}
+ else if (GOMP_MAP_NONCONTIG_ARRAY_P (kind & typemask))
+ {
+ /* Ignore non-contiguous arrays for now, we process them together
+ later. */
+ tgt->list[i].key = NULL;
+ tgt->list[i].offset = 0;
+ not_found_cnt++;
+
+ /* The map for the non-contiguous array itself is never copied from
+ during unmapping, its the data rows that count. Set copy-from
+ flags to false here. */
+ tgt->list[i].copy_from = false;
+ tgt->list[i].always_copy_from = false;
+ tgt->list[i].do_detach = false;
+
+ size_t align = (size_t) 1 << (kind >> rshift);
+ if (tgt_align < align)
+ tgt_align = align;
+
+ continue;
+ }
+
cur_node.host_start = (uintptr_t) hostaddrs[i];
if (!GOMP_MAP_POINTER_P (kind & typemask)
&& (kind & typemask) != GOMP_MAP_ATTACH)
@@ -898,6 +922,44 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
}
}
+ /* For non-contiguous arrays. Each data row is one target item, separated
+ from the normal map clause items, hence we order them after mapnum. */
+ if (nca_info)
+ {
+ struct target_var_desc *next_var_desc = &tgt->list[mapnum];
+ for (i = 0; i < nca_info->num_ncarray; i++)
+ {
+ struct goacc_ncarray *nca = &nca_info->ncarray[i];
+ int kind = get_kind (short_mapkind, kinds, nca->map_index);
+ size_t align = (size_t) 1 << (kind >> rshift);
+ tgt_size = (tgt_size + align - 1) & ~(align - 1);
+ tgt_size += nca->ptrblock_size;
+
+ for (size_t j = 0; j < nca->data_row_num; j++)
+ {
+ struct target_var_desc *row_desc = next_var_desc++;
+ void *row = nca->data_rows[j];
+ cur_node.host_start = (uintptr_t) row;
+ cur_node.host_end = cur_node.host_start + nca->data_row_size;
+ splay_tree_key n = splay_tree_lookup (mem_map, &cur_node);
+ if (n)
+ {
+ assert (n->refcount != REFCOUNT_LINK);
+ gomp_map_vars_existing (devicep, aq, n, &cur_node, row_desc,
+ kind & typemask,
+ /* TODO: cbuf? */ NULL);
+ }
+ else
+ {
+ tgt_size = (tgt_size + align - 1) & ~(align - 1);
+ tgt_size += nca->data_row_size;
+ not_found_cnt++;
+ }
+ }
+ }
+ assert (next_var_desc == &tgt->list[mapnum + nca_info->num_data_rows]);
+ }
+
if (devaddrs)
{
if (mapnum != 1)
@@ -1110,6 +1172,15 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
default:
break;
}
+
+ if (GOMP_MAP_NONCONTIG_ARRAY_P (kind & typemask))
+ {
+ tgt->list[i].key = &array->key;
+ tgt->list[i].key->tgt = tgt;
+ array++;
+ continue;
+ }
+
splay_tree_key k = &array->key;
k->host_start = (uintptr_t) hostaddrs[i];
if (!GOMP_MAP_POINTER_P (kind & typemask))
@@ -1263,6 +1334,99 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
array++;
}
}
+
+ /* Processing of non-contiguous array rows. */
+ if (nca_info)
+ {
+ struct target_var_desc *next_var_desc = &tgt->list[mapnum];
+ for (i = 0; i < nca_info->num_ncarray; i++)
+ {
+ struct goacc_ncarray *nca = &nca_info->ncarray[i];
+ int kind = get_kind (short_mapkind, kinds, nca->map_index);
+ size_t align = (size_t) 1 << (kind >> rshift);
+ tgt_size = (tgt_size + align - 1) & ~(align - 1);
+
+ assert (nca->ptr == hostaddrs[nca->map_index]);
+
+ /* For the map of the non-contiguous array itself, adjust so that
+ the passed device address points to the beginning of the
+ ptrblock. Remember to adjust the first-dimension's bias here. */
+ tgt->list[nca->map_index].key->tgt_offset
+ = tgt_size - nca->descr->dims[0].base;
+
+ void *target_ptrblock = (void*) tgt->tgt_start + tgt_size;
+ tgt_size += nca->ptrblock_size;
+
+ /* Add splay key for each data row in current non-contiguous
+ array. */
+ for (size_t j = 0; j < nca->data_row_num; j++)
+ {
+ struct target_var_desc *row_desc = next_var_desc++;
+ void *row = nca->data_rows[j];
+ cur_node.host_start = (uintptr_t) row;
+ cur_node.host_end = cur_node.host_start + nca->data_row_size;
+ splay_tree_key k = splay_tree_lookup (mem_map, &cur_node);
+ if (k)
+ {
+ assert (k->refcount != REFCOUNT_LINK);
+ gomp_map_vars_existing (devicep, aq, k, &cur_node, row_desc,
+ kind & typemask,
+ cbufp);
+ }
+ else
+ {
+ tgt->refcount++;
+ tgt_size = (tgt_size + align - 1) & ~(align - 1);
+
+ k = &array->key;
+ k->host_start = (uintptr_t) row;
+ k->host_end = k->host_start + nca->data_row_size;
+
+ k->tgt = tgt;
+ k->refcount = 1;
+ k->virtual_refcount = 0;
+ k->aux = NULL;
+ k->tgt_offset = tgt_size;
+
+ tgt_size += nca->data_row_size;
+
+ row_desc->key = k;
+ row_desc->copy_from
+ = GOMP_MAP_COPY_FROM_P (kind & typemask);
+ row_desc->always_copy_from
+ = GOMP_MAP_COPY_FROM_P (kind & typemask);
+ row_desc->do_detach = false;
+ row_desc->offset = 0;
+ row_desc->length = nca->data_row_size;
+
+ array->left = NULL;
+ array->right = NULL;
+ splay_tree_insert (mem_map, array);
+
+ if (GOMP_MAP_COPY_TO_P (kind & typemask))
+ gomp_copy_host2dev (devicep, aq,
+ (void *) tgt->tgt_start + k->tgt_offset,
+ (void *) k->host_start,
+ nca->data_row_size, cbufp);
+ array++;
+ }
+ nca->tgt_data_rows[j]
+ = (void *) (k->tgt->tgt_start + k->tgt_offset);
+ }
+
+ /* Now we have the target memory allocated, and target offsets of all
+ row blocks assigned and calculated, we can construct the
+ accelerator side ptrblock and copy it in. */
+ if (nca->ptrblock_size)
+ {
+ void *ptrblock = goacc_noncontig_array_create_ptrblock
+ (nca, target_ptrblock);
+ gomp_copy_host2dev (devicep, aq, target_ptrblock, ptrblock,
+ nca->ptrblock_size, cbufp);
+ free (ptrblock);
+ }
+ }
+ }
}
if (pragma_kind == GOMP_MAP_VARS_TARGET)
@@ -1316,13 +1480,26 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
return tgt;
}
+attribute_hidden struct target_mem_desc *
+gomp_map_vars_openacc (struct gomp_device_descr *devicep,
+ struct goacc_asyncqueue *aq, size_t mapnum,
+ void **hostaddrs, size_t *sizes, unsigned short *kinds,
+ void *nca_info)
+{
+ return gomp_map_vars_internal (devicep, aq, mapnum, hostaddrs, NULL,
+ sizes, (void *) kinds,
+ (struct goacc_ncarray_info *) nca_info,
+ true, GOMP_MAP_VARS_OPENACC);
+}
+
attribute_hidden struct target_mem_desc *
gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
void **hostaddrs, void **devaddrs, size_t *sizes, void *kinds,
bool short_mapkind, enum gomp_map_vars_kind pragma_kind)
{
return gomp_map_vars_internal (devicep, NULL, mapnum, hostaddrs, devaddrs,
- sizes, kinds, short_mapkind, pragma_kind);
+ sizes, kinds, NULL, short_mapkind,
+ pragma_kind);
}
attribute_hidden struct target_mem_desc *
@@ -1333,7 +1510,7 @@ gomp_map_vars_async (struct gomp_device_descr *devicep,
enum gomp_map_vars_kind pragma_kind)
{
return gomp_map_vars_internal (devicep, aq, mapnum, hostaddrs, devaddrs,
- sizes, kinds, short_mapkind, pragma_kind);
+ sizes, kinds, NULL, short_mapkind, pragma_kind);
}
static void
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-1.c
new file mode 100644
index 00000000000..a70375c03f4
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-1.c
@@ -0,0 +1,103 @@
+/* { dg-do run } */
+
+#include <stdlib.h>
+#include <assert.h>
+
+#define n 100
+#define m 100
+
+int b[n][m];
+
+void
+test1 (void)
+{
+ int i, j, *a[100];
+
+ /* Array of pointers form test. */
+ for (i = 0; i < n; i++)
+ {
+ a[i] = (int *)malloc (sizeof (int) * m);
+ for (j = 0; j < m; j++)
+ b[i][j] = j - i;
+ }
+
+ #pragma acc parallel loop copyout(a[0:n][0:m]) copyin(b)
+ for (i = 0; i < n; i++)
+ #pragma acc loop
+ for (j = 0; j < m; j++)
+ a[i][j] = b[i][j];
+
+ for (i = 0; i < n; i++)
+ {
+ for (j = 0; j < m; j++)
+ assert (a[i][j] == b[i][j]);
+ /* Clean up. */
+ free (a[i]);
+ }
+}
+
+void
+test2 (void)
+{
+ int i, j, **a = (int **) malloc (sizeof (int *) * n);
+
+ /* Separately allocated blocks. */
+ for (i = 0; i < n; i++)
+ {
+ a[i] = (int *)malloc (sizeof (int) * m);
+ for (j = 0; j < m; j++)
+ b[i][j] = j - i;
+ }
+
+ #pragma acc parallel loop copyout(a[0:n][0:m]) copyin(b)
+ for (i = 0; i < n; i++)
+ #pragma acc loop
+ for (j = 0; j < m; j++)
+ a[i][j] = b[i][j];
+
+ for (i = 0; i < n; i++)
+ {
+ for (j = 0; j < m; j++)
+ assert (a[i][j] == b[i][j]);
+ /* Clean up. */
+ free (a[i]);
+ }
+ free (a);
+}
+
+void
+test3 (void)
+{
+ int i, j, **a = (int **) malloc (sizeof (int *) * n);
+ a[0] = (int *) malloc (sizeof (int) * n * m);
+
+ /* Rows allocated in one contiguous block. */
+ for (i = 0; i < n; i++)
+ {
+ a[i] = *a + i * m;
+ for (j = 0; j < m; j++)
+ b[i][j] = j - i;
+ }
+
+ #pragma acc parallel loop copyout(a[0:n][0:m]) copyin(b)
+ for (i = 0; i < n; i++)
+ #pragma acc loop
+ for (j = 0; j < m; j++)
+ a[i][j] = b[i][j];
+
+ for (i = 0; i < n; i++)
+ for (j = 0; j < m; j++)
+ assert (a[i][j] == b[i][j]);
+
+ free (a[0]);
+ free (a);
+}
+
+int
+main (void)
+{
+ test1 ();
+ test2 ();
+ test3 ();
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-2.c
new file mode 100644
index 00000000000..b85c6371f25
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-2.c
@@ -0,0 +1,37 @@
+/* { dg-do run } */
+
+#include <assert.h>
+#include "noncontig_array-utils.h"
+
+int
+main (void)
+{
+ int n = 10;
+ int ***a = (int ***) create_ncarray (sizeof (int), n, 3);
+ int ***b = (int ***) create_ncarray (sizeof (int), n, 3);
+ int ***c = (int ***) create_ncarray (sizeof (int), n, 3);
+
+ for (int i = 0; i < n; i++)
+ for (int j = 0; j < n; j++)
+ for (int k = 0; k < n; k++)
+ {
+ a[i][j][k] = i + j * k + k;
+ b[i][j][k] = j + k * i + i * j;
+ c[i][j][k] = a[i][j][k];
+ }
+
+ #pragma acc parallel copy (a[0:n][0:n][0:n]) copyin (b[0:n][0:n][0:n])
+ {
+ for (int i = 0; i < n; i++)
+ for (int j = 0; j < n; j++)
+ for (int k = 0; k < n; k++)
+ a[i][j][k] += b[k][j][i] + i + j + k;
+ }
+
+ for (int i = 0; i < n; i++)
+ for (int j = 0; j < n; j++)
+ for (int k = 0; k < n; k++)
+ assert (a[i][j][k] == c[i][j][k] + b[k][j][i] + i + j + k);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-3.c
new file mode 100644
index 00000000000..99db207493e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-3.c
@@ -0,0 +1,45 @@
+/* { dg-do run } */
+
+#include <assert.h>
+#include "noncontig_array-utils.h"
+
+int main (void)
+{
+ int n = 20, x = 5, y = 12;
+ int *****a = (int *****) create_ncarray (sizeof (int), n, 5);
+
+ int sum1 = 0, sum2 = 0, sum3 = 0;
+
+ for (int i = 0; i < n; i++)
+ for (int j = 0; j < n; j++)
+ for (int k = 0; k < n; k++)
+ for (int l = 0; l < n; l++)
+ for (int m = 0; m < n; m++)
+ {
+ a[i][j][k][l][m] = 1;
+ sum1++;
+ }
+
+ #pragma acc parallel copy (a[x:y][x:y][x:y][x:y][x:y]) copy(sum2)
+ {
+ for (int i = x; i < x + y; i++)
+ for (int j = x; j < x + y; j++)
+ for (int k = x; k < x + y; k++)
+ for (int l = x; l < x + y; l++)
+ for (int m = x; m < x + y; m++)
+ {
+ a[i][j][k][l][m] = 0;
+ sum2++;
+ }
+ }
+
+ for (int i = 0; i < n; i++)
+ for (int j = 0; j < n; j++)
+ for (int k = 0; k < n; k++)
+ for (int l = 0; l < n; l++)
+ for (int m = 0; m < n; m++)
+ sum3 += a[i][j][k][l][m];
+
+ assert (sum1 == sum2 + sum3);
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-4.c
new file mode 100644
index 00000000000..6cfaf98d37e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-4.c
@@ -0,0 +1,36 @@
+/* { dg-do run } */
+
+#include <assert.h>
+#include "noncontig_array-utils.h"
+
+int main (void)
+{
+ int n = 128;
+ double ***a = (double ***) create_ncarray (sizeof (double), n, 3);
+ double ***b = (double ***) create_ncarray (sizeof (double), n, 3);
+
+ for (int i = 0; i < n; i++)
+ for (int j = 0; j < n; j++)
+ for (int k = 0; k < n; k++)
+ a[i][j][k] = i + j + k + i * j * k;
+
+ /* This test exercises async copyout of non-contiguous array rows. */
+ #pragma acc parallel copyin(a[0:n][0:n][0:n]) copyout(b[0:n][0:n][0:n]) async(5)
+ {
+ #pragma acc loop gang
+ for (int i = 0; i < n; i++)
+ #pragma acc loop vector
+ for (int j = 0; j < n; j++)
+ for (int k = 0; k < n; k++)
+ b[i][j][k] = a[i][j][k] * 2.0;
+ }
+
+ #pragma acc wait (5)
+
+ for (int i = 0; i < n; i++)
+ for (int j = 0; j < n; j++)
+ for (int k = 0; k < n; k++)
+ assert (b[i][j][k] == a[i][j][k] * 2.0);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-utils.h b/libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-utils.h
new file mode 100644
index 00000000000..554bda77bbd
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-utils.h
@@ -0,0 +1,44 @@
+#include <stdlib.h>
+#include <string.h>
+#include <assert.h>
+#include <stdint.h>
+
+/* Allocate and create a pointer based NDIMS-dimensional array,
+ each dimension DIMLEN long, with ELSIZE sized data elements. */
+void *
+create_ncarray (size_t elsize, int dimlen, int ndims)
+{
+ size_t blk_size = 0;
+ size_t n = 1;
+
+ for (int i = 0; i < ndims - 1; i++)
+ {
+ n *= dimlen;
+ blk_size += sizeof (void *) * n;
+ }
+ size_t data_rows_num = n;
+ size_t data_rows_offset = blk_size;
+ blk_size += elsize * n * dimlen;
+
+ void *blk = (void *) malloc (blk_size);
+ memset (blk, 0, blk_size);
+ void **curr_dim = (void **) blk;
+ n = 1;
+
+ for (int d = 0; d < ndims - 1; d++)
+ {
+ uintptr_t next_dim = (uintptr_t) (curr_dim + n * dimlen);
+ size_t next_dimlen = dimlen * (d < ndims - 2 ? sizeof (void *) : elsize);
+
+ for (int b = 0; b < n; b++)
+ for (int i = 0; i < dimlen; i++)
+ if (d < ndims - 1)
+ curr_dim[b * dimlen + i]
+ = (void*) (next_dim + b * dimlen * next_dimlen + i * next_dimlen);
+
+ n *= dimlen;
+ curr_dim = (void**) next_dim;
+ }
+ assert (n == data_rows_num);
+ return blk;
+}
^ permalink raw reply [flat|nested] 6+ messages in thread
* Re: [PATCH, OpenACC, v3] Non-contiguous array support for OpenACC data clauses
2019-11-26 14:57 [PATCH, OpenACC, v3] Non-contiguous array support for OpenACC data clauses Chung-Lin Tang
2020-04-07 13:50 ` (v4 update) " Chung-Lin Tang
@ 2023-03-10 13:24 ` Thomas Schwinge
2023-03-15 14:47 ` Thomas Schwinge
2023-04-03 14:39 ` [og12] OpenACC: Pass pre-allocated 'ptrblock' to 'goacc_noncontig_array_create_ptrblock' [PR76739] (was: [PATCH, OpenACC, v3] Non-contiguous array support for OpenACC data clauses) Thomas Schwinge
3 siblings, 0 replies; 6+ messages in thread
From: Thomas Schwinge @ 2023-03-10 13:24 UTC (permalink / raw)
To: Chung-Lin Tang, gcc-patches
[-- Attachment #1: Type: text/plain, Size: 2587 bytes --]
Hi!
On 2019-11-26T22:49:21+0800, Chung-Lin Tang <chunglin_tang@mentor.com> wrote:
> this is a reorg of the last non-contiguous arrays patch.
(Sorry, this is still not the master branch integration email...)
Just a small clean-up, to simplify other changes that I'm working on:
> (4) Along the way, I've added a 'gomp_map_vars_openacc' for specializing our
> uses, which should shave off quite some code through inlining.
> --- libgomp/libgomp.h (revision 278656)
> +++ libgomp/libgomp.h (working copy)
> @@ -1167,6 +1167,10 @@ extern struct target_mem_desc *gomp_map_vars_async
> size_t, void **, void **,
> size_t *, void *, bool,
> enum gomp_map_vars_kind);
> +extern struct target_mem_desc *gomp_map_vars_openacc (struct gomp_device_descr *,
> + struct goacc_asyncqueue *,
> + size_t, void **, size_t *,
> + unsigned short *, void *);
> extern void gomp_unmap_tgt (struct target_mem_desc *);
> extern void gomp_unmap_vars (struct target_mem_desc *, bool);
> extern void gomp_unmap_vars_async (struct target_mem_desc *, bool,
> --- libgomp/target.c (revision 278656)
> +++ libgomp/target.c (working copy)
> @@ -1086,12 +1248,25 @@ gomp_map_vars_internal (struct gomp_device_descr *
> }
>
> attribute_hidden struct target_mem_desc *
> +gomp_map_vars_openacc (struct gomp_device_descr *devicep,
> + struct goacc_asyncqueue *aq, size_t mapnum,
> + void **hostaddrs, size_t *sizes, unsigned short *kinds,
> + void *nca_info)
> +{
> + return gomp_map_vars_internal (devicep, aq, mapnum, hostaddrs, NULL,
> + sizes, (void *) kinds,
> + (struct goacc_ncarray_info *) nca_info,
> + true, GOMP_MAP_VARS_OPENACC);
> +}
Pushed to devel/omp/gcc-12 branch
commit 5ea330fdc918e6731c5b706715a18470909247bf
"libgomp: Merge 'gomp_map_vars_openacc' into 'goacc_map_vars' [PR76739]",
see attached.
Grüße
Thomas
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-libgomp-Merge-gomp_map_vars_openacc-into-goacc_map_v.patch --]
[-- Type: text/x-diff, Size: 7495 bytes --]
From 5ea330fdc918e6731c5b706715a18470909247bf Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Thu, 2 Mar 2023 18:36:47 +0100
Subject: [PATCH] libgomp: Merge 'gomp_map_vars_openacc' into 'goacc_map_vars'
[PR76739]
Upstream has 'goacc_map_vars'; merge the new 'gomp_map_vars_openacc' into it.
(Maybe the latter didn't exist yet when the former was originally added?)
No functional change.
Clean-up for og12 commit 15d0f61a7fecdc8fd12857c40879ea3730f6d99f
"Merge non-contiguous array support patches".
PR other/76739
libgomp/
* libgomp.h (goacc_map_vars): Add 'struct goacc_ncarray_info *'
formal parameter.
(gomp_map_vars_openacc): Remove.
* target.c (goacc_map_vars): Adjust.
(gomp_map_vars_openacc): Remove.
* oacc-mem.c (acc_map_data, goacc_enter_datum)
(goacc_enter_data_internal): Adjust.
* oacc-parallel.c (GOACC_parallel_keyed, GOACC_data_start):
Adjust.
---
libgomp/ChangeLog.omp | 11 +++++++++++
libgomp/libgomp.h | 9 ++++-----
libgomp/oacc-mem.c | 8 ++++----
libgomp/oacc-parallel.c | 10 +++++-----
libgomp/target.c | 17 +++--------------
5 files changed, 27 insertions(+), 28 deletions(-)
diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index 0e984754bb0..be21ec39428 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,5 +1,16 @@
2023-03-10 Thomas Schwinge <thomas@codesourcery.com>
+ PR other/76739
+ * libgomp.h (goacc_map_vars): Add 'struct goacc_ncarray_info *'
+ formal parameter.
+ (gomp_map_vars_openacc): Remove.
+ * target.c (goacc_map_vars): Adjust.
+ (gomp_map_vars_openacc): Remove.
+ * oacc-mem.c (acc_map_data, goacc_enter_datum)
+ (goacc_enter_data_internal): Adjust.
+ * oacc-parallel.c (GOACC_parallel_keyed, GOACC_data_start):
+ Adjust.
+
* oacc-host.c: Revert
"OpenACC profiling-interface fixes for asynchronous operations"
changes.
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index ba12d558465..92f6f14960f 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -1445,15 +1445,14 @@ extern void gomp_attach_pointer (struct gomp_device_descr *,
extern void gomp_detach_pointer (struct gomp_device_descr *,
struct goacc_asyncqueue *, splay_tree_key,
uintptr_t, bool, struct gomp_coalesce_buf *);
+struct goacc_ncarray_info;
extern struct target_mem_desc *goacc_map_vars (struct gomp_device_descr *,
struct goacc_asyncqueue *,
size_t, void **, void **,
- size_t *, void *, bool,
+ size_t *, void *,
+ struct goacc_ncarray_info *,
+ bool,
enum gomp_map_vars_kind);
-extern struct target_mem_desc *gomp_map_vars_openacc (struct gomp_device_descr *,
- struct goacc_asyncqueue *,
- size_t, void **, size_t *,
- unsigned short *, void *);
extern void goacc_unmap_vars (struct target_mem_desc *, bool,
struct goacc_asyncqueue *);
extern void gomp_init_device (struct gomp_device_descr *);
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index 6fb8be98542..bd82beefcdb 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -403,7 +403,7 @@ acc_map_data (void *h, void *d, size_t s)
struct target_mem_desc *tgt
= goacc_map_vars (acc_dev, NULL, mapnum, &hostaddrs, &devaddrs, &sizes,
- &kinds, true, GOMP_MAP_VARS_ENTER_DATA);
+ &kinds, NULL, true, GOMP_MAP_VARS_ENTER_DATA);
assert (tgt);
assert (tgt->list_count == 1);
splay_tree_key n = tgt->list[0].key;
@@ -572,7 +572,7 @@ goacc_enter_datum (void **hostaddrs, size_t *sizes, void *kinds, int async)
struct target_mem_desc *tgt
= goacc_map_vars (acc_dev, aq, mapnum, hostaddrs, NULL, sizes,
- kinds, true, GOMP_MAP_VARS_ENTER_DATA);
+ kinds, NULL, true, GOMP_MAP_VARS_ENTER_DATA);
assert (tgt);
assert (tgt->list_count == 1);
n = tgt->list[0].key;
@@ -1247,7 +1247,7 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
gomp_mutex_unlock (&acc_dev->lock);
struct target_mem_desc *tgt_
= goacc_map_vars (acc_dev, aq, groupnum, &hostaddrs[i], NULL,
- &sizes[i], &kinds[i], true,
+ &sizes[i], &kinds[i], NULL, true,
GOMP_MAP_VARS_ENTER_DATA);
assert (tgt_ == NULL);
gomp_mutex_lock (&acc_dev->lock);
@@ -1297,7 +1297,7 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
struct target_mem_desc *tgt
= goacc_map_vars (acc_dev, aq, groupnum, &hostaddrs[i], NULL,
- &sizes[i], &kinds[i], true,
+ &sizes[i], &kinds[i], NULL, true,
GOMP_MAP_VARS_ENTER_DATA);
assert (tgt);
diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c
index d66bc882a5f..9c1db402c82 100644
--- a/libgomp/oacc-parallel.c
+++ b/libgomp/oacc-parallel.c
@@ -468,8 +468,8 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (void *),
goacc_aq aq = get_goacc_asyncqueue (async);
- tgt = gomp_map_vars_openacc (acc_dev, aq, mapnum, hostaddrs, sizes, kinds,
- nca_info);
+ tgt = goacc_map_vars (acc_dev, aq, mapnum, hostaddrs, NULL, sizes, kinds,
+ nca_info, true, 0);
free (nca_info);
if (profiling_p)
@@ -635,7 +635,7 @@ GOACC_data_start (int flags_m, size_t mapnum,
{
prof_info.device_type = acc_device_host;
api_info.device_type = prof_info.device_type;
- tgt = gomp_map_vars_openacc (NULL, NULL, 0, NULL, NULL, NULL, NULL);
+ tgt = goacc_map_vars (NULL, NULL, 0, NULL, NULL, NULL, NULL, NULL, true, 0);
tgt->prev = thr->mapped_data;
thr->mapped_data = tgt;
@@ -652,8 +652,8 @@ GOACC_data_start (int flags_m, size_t mapnum,
}
gomp_debug (0, " %s: prepare mappings\n", __FUNCTION__);
- tgt = gomp_map_vars_openacc (acc_dev, NULL, mapnum, hostaddrs, sizes, kinds,
- nca_info);
+ tgt = goacc_map_vars (acc_dev, NULL, mapnum, hostaddrs, NULL, sizes, kinds,
+ nca_info, true, 0);
free (nca_info);
gomp_debug (0, " %s: mappings prepared\n", __FUNCTION__);
tgt->prev = thr->mapped_data;
diff --git a/libgomp/target.c b/libgomp/target.c
index fcc5b9dabca..e4fc7da6f07 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -1986,18 +1986,6 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
return tgt;
}
-attribute_hidden struct target_mem_desc *
-gomp_map_vars_openacc (struct gomp_device_descr *devicep,
- struct goacc_asyncqueue *aq, size_t mapnum,
- void **hostaddrs, size_t *sizes, unsigned short *kinds,
- void *nca_info)
-{
- return gomp_map_vars_internal (devicep, aq, mapnum, hostaddrs, NULL,
- sizes, (void *) kinds,
- (struct goacc_ncarray_info *) nca_info,
- true, NULL, GOMP_MAP_VARS_OPENACC);
-}
-
static struct target_mem_desc *
gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
void **hostaddrs, void **devaddrs, size_t *sizes, void *kinds,
@@ -2027,11 +2015,12 @@ attribute_hidden struct target_mem_desc *
goacc_map_vars (struct gomp_device_descr *devicep,
struct goacc_asyncqueue *aq, size_t mapnum,
void **hostaddrs, void **devaddrs, size_t *sizes,
- void *kinds, bool short_mapkind,
+ void *kinds, struct goacc_ncarray_info *nca_info,
+ bool short_mapkind,
enum gomp_map_vars_kind pragma_kind)
{
return gomp_map_vars_internal (devicep, aq, mapnum, hostaddrs, devaddrs,
- sizes, kinds, NULL, short_mapkind, NULL,
+ sizes, kinds, nca_info, short_mapkind, NULL,
GOMP_MAP_VARS_OPENACC | pragma_kind);
}
--
2.25.1
^ permalink raw reply [flat|nested] 6+ messages in thread
* Re: [PATCH, OpenACC, v3] Non-contiguous array support for OpenACC data clauses
2019-11-26 14:57 [PATCH, OpenACC, v3] Non-contiguous array support for OpenACC data clauses Chung-Lin Tang
2020-04-07 13:50 ` (v4 update) " Chung-Lin Tang
2023-03-10 13:24 ` Thomas Schwinge
@ 2023-03-15 14:47 ` Thomas Schwinge
2023-03-24 15:17 ` Thomas Schwinge
2023-04-03 14:39 ` [og12] OpenACC: Pass pre-allocated 'ptrblock' to 'goacc_noncontig_array_create_ptrblock' [PR76739] (was: [PATCH, OpenACC, v3] Non-contiguous array support for OpenACC data clauses) Thomas Schwinge
3 siblings, 1 reply; 6+ messages in thread
From: Thomas Schwinge @ 2023-03-15 14:47 UTC (permalink / raw)
To: Chung-Lin Tang; +Cc: gcc-patches
[-- Attachment #1: Type: text/plain, Size: 2542 bytes --]
Hi Chung-Lin!
On 2019-11-26T22:49:21+0800, Chung-Lin Tang <chunglin_tang@mentor.com> wrote:
> this is a reorg of the last non-contiguous arrays patch.
(Sorry, this is still not the master branch integration email...)
I noticed the following while working on something else:
> --- libgomp/oacc-parallel.c (revision 278656)
> +++ libgomp/oacc-parallel.c (working copy)
> @@ -311,8 +488,10 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (voi
>
> goacc_aq aq = get_goacc_asyncqueue (async);
>
> - tgt = gomp_map_vars_async (acc_dev, aq, mapnum, hostaddrs, NULL, sizes, kinds,
> - true, GOMP_MAP_VARS_OPENACC);
> + tgt = gomp_map_vars_openacc (acc_dev, aq, mapnum, hostaddrs, sizes, kinds,
> + nca_info);
> + free (nca_info);
Given OpenACC 'async', don't we have to defer 'free' of the
non-contiguous array support data structure here? But I'm not completely
sure -- can we rule out that any asynchronous copying of any data of
'nca_info' is still going on after returning from 'gomp_map_vars'?
> @@ -488,9 +666,19 @@ GOACC_data_start (int flags_m, size_t mapnum,
> - tgt = gomp_map_vars (acc_dev, mapnum, hostaddrs, NULL, sizes, kinds, true,
> - GOMP_MAP_VARS_OPENACC);
> + tgt = gomp_map_vars_openacc (acc_dev, NULL, mapnum, hostaddrs, sizes, kinds,
> + nca_info);
> + free (nca_info);
Here, it's not relevant, as there is no 'async' support (yet) for 'data'
constructs.
> --- libgomp/target.c (revision 278656)
> +++ libgomp/target.c (working copy)
> @@ -1044,6 +1114,98 @@ gomp_map_vars_internal (struct gomp_device_descr *
> + void *ptrblock = goacc_noncontig_array_create_ptrblock
> + (nca, target_ptrblock);
> + gomp_copy_host2dev (devicep, aq, target_ptrblock, ptrblock,
> + nca->ptrblock_size, cbufp);
> + free (ptrblock);
Here again, however, don't we have to defer the 'free'?
Please verify the attached
"Given OpenACC 'async', defer 'free' of non-contiguous array support data structures",
in particular the 'libgomp/oacc-parallel.c:GOACC_parallel_keyed' case.
Grüße
Thomas
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-Given-OpenACC-async-defer-free-of-non-contiguous-arr.patch --]
[-- Type: text/x-diff, Size: 1794 bytes --]
From 998f1156a51010490c2e918defb4517706916c92 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Wed, 15 Mar 2023 13:34:02 +0100
Subject: [PATCH] Given OpenACC 'async', defer 'free' of non-contiguous array
support data structures
Fix-up for og12 commit 15d0f61a7fecdc8fd12857c40879ea3730f6d99f
"Merge non-contiguous array support patches".
libgomp/
* oacc-parallel.c (GOACC_parallel_keyed): Given OpenACC 'async',
defer 'free' of non-contiguous array support data structures.
* target.c (gomp_map_vars_internal): Likewise.
---
libgomp/oacc-parallel.c | 5 ++++-
libgomp/target.c | 6 +++++-
2 files changed, 9 insertions(+), 2 deletions(-)
diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c
index 9cd99b4a0b4..136702d6e61 100644
--- a/libgomp/oacc-parallel.c
+++ b/libgomp/oacc-parallel.c
@@ -470,7 +470,10 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (void *),
= goacc_map_vars (acc_dev, aq, mapnum, hostaddrs, NULL, sizes, kinds,
nca_info, true, GOMP_MAP_VARS_TARGET);
- free (nca_info);
+ if (aq == NULL)
+ free (nca_info);
+ else
+ acc_dev->openacc.async.queue_callback_func (aq, free, nca_info);
if (profiling_p)
{
diff --git a/libgomp/target.c b/libgomp/target.c
index 96ece0b31fd..aaa597f6610 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -1938,7 +1938,11 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
(nca, target_ptrblock);
gomp_copy_host2dev (devicep, aq, target_ptrblock, ptrblock,
nca->ptrblock_size, false, cbufp);
- free (ptrblock);
+ if (aq)
+ /* Free once the transfer has completed. */
+ devicep->openacc.async.queue_callback_func (aq, free, ptrblock);
+ else
+ free (ptrblock);
}
}
}
--
2.25.1
^ permalink raw reply [flat|nested] 6+ messages in thread
* Re: [PATCH, OpenACC, v3] Non-contiguous array support for OpenACC data clauses
2023-03-15 14:47 ` Thomas Schwinge
@ 2023-03-24 15:17 ` Thomas Schwinge
0 siblings, 0 replies; 6+ messages in thread
From: Thomas Schwinge @ 2023-03-24 15:17 UTC (permalink / raw)
To: Chung-Lin Tang, gcc-patches
[-- Attachment #1: Type: text/plain, Size: 2915 bytes --]
Hi!
On 2023-03-15T15:47:47+0100, I wrote:
> On 2019-11-26T22:49:21+0800, Chung-Lin Tang <chunglin_tang@mentor.com> wrote:
>> this is a reorg of the last non-contiguous arrays patch.
>
> (Sorry, this is still not the master branch integration email...)
>
>
> I noticed the following while working on something else:
>
>> --- libgomp/oacc-parallel.c (revision 278656)
>> +++ libgomp/oacc-parallel.c (working copy)
>
>> @@ -311,8 +488,10 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (voi
>>
>> goacc_aq aq = get_goacc_asyncqueue (async);
>>
>> - tgt = gomp_map_vars_async (acc_dev, aq, mapnum, hostaddrs, NULL, sizes, kinds,
>> - true, GOMP_MAP_VARS_OPENACC);
>> + tgt = gomp_map_vars_openacc (acc_dev, aq, mapnum, hostaddrs, sizes, kinds,
>> + nca_info);
>> + free (nca_info);
>
> Given OpenACC 'async', don't we have to defer 'free' of the
> non-contiguous array support data structure here? But I'm not completely
> sure -- can we rule out that any asynchronous copying of any data of
> 'nca_info' is still going on after returning from 'gomp_map_vars'?
>
>> @@ -488,9 +666,19 @@ GOACC_data_start (int flags_m, size_t mapnum,
>
>> - tgt = gomp_map_vars (acc_dev, mapnum, hostaddrs, NULL, sizes, kinds, true,
>> - GOMP_MAP_VARS_OPENACC);
>> + tgt = gomp_map_vars_openacc (acc_dev, NULL, mapnum, hostaddrs, sizes, kinds,
>> + nca_info);
>> + free (nca_info);
>
> Here, it's not relevant, as there is no 'async' support (yet) for 'data'
> constructs.
>
>> --- libgomp/target.c (revision 278656)
>> +++ libgomp/target.c (working copy)
>
>> @@ -1044,6 +1114,98 @@ gomp_map_vars_internal (struct gomp_device_descr *
>
>> + void *ptrblock = goacc_noncontig_array_create_ptrblock
>> + (nca, target_ptrblock);
>> + gomp_copy_host2dev (devicep, aq, target_ptrblock, ptrblock,
>> + nca->ptrblock_size, cbufp);
>> + free (ptrblock);
>
> Here again, however, don't we have to defer the 'free'?
>
> Please verify the attached
> "Given OpenACC 'async', defer 'free' of non-contiguous array support data structures",
> in particular the 'libgomp/oacc-parallel.c:GOACC_parallel_keyed' case.
To allow me to make progress with the "something else" (depending on
this), I've now pushed to devel/omp/gcc-12
commit a1f6758ae08fa748b291954371859e0158d4d667
"Given OpenACC 'async', defer 'free' of non-contiguous array support data structures [PR76739]",
see attached.
Grüße
Thomas
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-Given-OpenACC-async-defer-free-of-non-contiguous-arr.patch --]
[-- Type: text/x-diff, Size: 2403 bytes --]
From a1f6758ae08fa748b291954371859e0158d4d667 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Wed, 15 Mar 2023 13:34:02 +0100
Subject: [PATCH] Given OpenACC 'async', defer 'free' of non-contiguous array
support data structures [PR76739]
Fix-up for og12 commit 15d0f61a7fecdc8fd12857c40879ea3730f6d99f
"Merge non-contiguous array support patches".
PR other/76739
libgomp/
* oacc-parallel.c (GOACC_parallel_keyed): Given OpenACC 'async',
defer 'free' of non-contiguous array support data structures.
* target.c (gomp_map_vars_internal): Likewise.
---
libgomp/ChangeLog.omp | 7 +++++++
libgomp/oacc-parallel.c | 5 ++++-
libgomp/target.c | 6 +++++-
3 files changed, 16 insertions(+), 2 deletions(-)
diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index ace54f2f82f..85ebab14ba8 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,10 @@
+2023-03-24 Thomas Schwinge <thomas@codesourcery.com>
+
+ PR other/76739
+ * oacc-parallel.c (GOACC_parallel_keyed): Given OpenACC 'async',
+ defer 'free' of non-contiguous array support data structures.
+ * target.c (gomp_map_vars_internal): Likewise.
+
2023-03-23 Tobias Burnus <tobias@codesourcery.com>
* testsuite/libgomp.fortran/map-alloc-comp-8.f90: New test.
diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c
index 9cd99b4a0b4..136702d6e61 100644
--- a/libgomp/oacc-parallel.c
+++ b/libgomp/oacc-parallel.c
@@ -470,7 +470,10 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (void *),
= goacc_map_vars (acc_dev, aq, mapnum, hostaddrs, NULL, sizes, kinds,
nca_info, true, GOMP_MAP_VARS_TARGET);
- free (nca_info);
+ if (aq == NULL)
+ free (nca_info);
+ else
+ acc_dev->openacc.async.queue_callback_func (aq, free, nca_info);
if (profiling_p)
{
diff --git a/libgomp/target.c b/libgomp/target.c
index 96ece0b31fd..aaa597f6610 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -1938,7 +1938,11 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
(nca, target_ptrblock);
gomp_copy_host2dev (devicep, aq, target_ptrblock, ptrblock,
nca->ptrblock_size, false, cbufp);
- free (ptrblock);
+ if (aq)
+ /* Free once the transfer has completed. */
+ devicep->openacc.async.queue_callback_func (aq, free, ptrblock);
+ else
+ free (ptrblock);
}
}
}
--
2.25.1
^ permalink raw reply [flat|nested] 6+ messages in thread
* [og12] OpenACC: Pass pre-allocated 'ptrblock' to 'goacc_noncontig_array_create_ptrblock' [PR76739] (was: [PATCH, OpenACC, v3] Non-contiguous array support for OpenACC data clauses)
2019-11-26 14:57 [PATCH, OpenACC, v3] Non-contiguous array support for OpenACC data clauses Chung-Lin Tang
` (2 preceding siblings ...)
2023-03-15 14:47 ` Thomas Schwinge
@ 2023-04-03 14:39 ` Thomas Schwinge
3 siblings, 0 replies; 6+ messages in thread
From: Thomas Schwinge @ 2023-04-03 14:39 UTC (permalink / raw)
To: Chung-Lin Tang, gcc-patches
[-- Attachment #1: Type: text/plain, Size: 1954 bytes --]
Hi!
On 2019-11-26T22:49:21+0800, Chung-Lin Tang <chunglin_tang@mentor.com> wrote:
> this is a reorg of the last non-contiguous arrays patch.
(Sorry, this is still not the master branch integration email...)
Just a small clean-up, to simplify other changes that I'm working on:
On 2019-11-26T22:49:21+0800, Chung-Lin Tang <chunglin_tang@mentor.com> wrote:
> --- libgomp/oacc-parallel.c (revision 278656)
> +++ libgomp/oacc-parallel.c (working copy)
> +void *
> +goacc_noncontig_array_create_ptrblock (struct goacc_ncarray *nca,
> + void *tgt_ptrblock_addr)
> +{
> + [...]
> + void *ptrblock = gomp_malloc (nca->ptrblock_size);
> --- libgomp/target.c (revision 278656)
> +++ libgomp/target.c (working copy)
> @@ -1044,6 +1114,98 @@ gomp_map_vars_internal (struct gomp_device_descr *
> + /* Now we have the target memory allocated, and target offsets of all
> + row blocks assigned and calculated, we can construct the
> + accelerator side ptrblock and copy it in. */
> + if (nca->ptrblock_size)
> + {
> + void *ptrblock = goacc_noncontig_array_create_ptrblock
> + (nca, target_ptrblock);
> + gomp_copy_host2dev (devicep, aq, target_ptrblock, ptrblock,
> + nca->ptrblock_size, cbufp);
> + free (ptrblock);
> + }
Pushed to devel/omp/gcc-12 branch
commit c58b28cb650995a41e1ab0166169799f3991bdd6
"OpenACC: Pass pre-allocated 'ptrblock' to 'goacc_noncontig_array_create_ptrblock' [PR76739]",
see attached.
Grüße
Thomas
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-OpenACC-Pass-pre-allocated-ptrblock-to-goacc_noncont.patch --]
[-- Type: text/x-diff, Size: 3500 bytes --]
From c58b28cb650995a41e1ab0166169799f3991bdd6 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Wed, 15 Mar 2023 14:32:12 +0100
Subject: [PATCH] OpenACC: Pass pre-allocated 'ptrblock' to
'goacc_noncontig_array_create_ptrblock' [PR76739]
... to simplify later changes. No functional change.
Follow-up for og12 commit 15d0f61a7fecdc8fd12857c40879ea3730f6d99f
"Merge non-contiguous array support patches".
PR other/76739
libgomp/
* target.c (gomp_map_vars_internal): Pass pre-allocated 'ptrblock'
to 'goacc_noncontig_array_create_ptrblock'.
* oacc-parallel.c (goacc_noncontig_array_create_ptrblock): Adjust.
* oacc-int.h (goacc_noncontig_array_create_ptrblock): Adjust.
---
libgomp/ChangeLog.omp | 6 ++++++
libgomp/oacc-int.h | 3 ++-
libgomp/oacc-parallel.c | 5 ++---
libgomp/target.c | 5 +++--
4 files changed, 13 insertions(+), 6 deletions(-)
diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index d8a7e476090..7afb5f43c04 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,5 +1,11 @@
2023-04-03 Thomas Schwinge <thomas@codesourcery.com>
+ PR other/76739
+ * target.c (gomp_map_vars_internal): Pass pre-allocated 'ptrblock'
+ to 'goacc_noncontig_array_create_ptrblock'.
+ * oacc-parallel.c (goacc_noncontig_array_create_ptrblock): Adjust.
+ * oacc-int.h (goacc_noncontig_array_create_ptrblock): Adjust.
+
* libgomp.texi (AMD Radeon, nvptx): Document OpenMP 'pinned'
memory.
diff --git a/libgomp/oacc-int.h b/libgomp/oacc-int.h
index d86aeb82dfa..28a6118873a 100644
--- a/libgomp/oacc-int.h
+++ b/libgomp/oacc-int.h
@@ -213,7 +213,8 @@ struct goacc_ncarray_info
struct goacc_ncarray ncarray[];
};
-extern void *goacc_noncontig_array_create_ptrblock (struct goacc_ncarray *, void *);
+extern void goacc_noncontig_array_create_ptrblock (struct goacc_ncarray *,
+ void *, void *);
#ifdef HAVE_ATTRIBUTE_VISIBILITY
diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c
index 136702d6e61..8d1c2cce836 100644
--- a/libgomp/oacc-parallel.c
+++ b/libgomp/oacc-parallel.c
@@ -165,13 +165,13 @@ goacc_process_noncontiguous_arrays (size_t mapnum, void **hostaddrs,
return nca_info;
}
-void *
+void
goacc_noncontig_array_create_ptrblock (struct goacc_ncarray *nca,
+ void *ptrblock,
void *tgt_ptrblock_addr)
{
struct goacc_ncarray_descr_type *descr = nca->descr;
void **tgt_data_rows = nca->tgt_data_rows;
- void *ptrblock = gomp_malloc (nca->ptrblock_size);
void **curr_dim_ptrblock = (void **) ptrblock;
size_t n = 1;
@@ -210,7 +210,6 @@ goacc_noncontig_array_create_ptrblock (struct goacc_ncarray *nca,
curr_dim_ptrblock = next_dim_ptrblock;
}
assert (n == nca->data_row_num);
- return ptrblock;
}
/* Handle the mapping pair that are presented when a
diff --git a/libgomp/target.c b/libgomp/target.c
index de3facb6428..b88b1ebaa13 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -1939,8 +1939,9 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
accelerator side ptrblock and copy it in. */
if (nca->ptrblock_size)
{
- void *ptrblock = goacc_noncontig_array_create_ptrblock
- (nca, target_ptrblock);
+ void *ptrblock = gomp_malloc (nca->ptrblock_size);
+ goacc_noncontig_array_create_ptrblock
+ (nca, ptrblock, target_ptrblock);
gomp_copy_host2dev (devicep, aq, target_ptrblock, ptrblock,
nca->ptrblock_size, false, cbufp);
if (aq)
--
2.25.1
^ permalink raw reply [flat|nested] 6+ messages in thread
end of thread, other threads:[~2023-04-03 14:39 UTC | newest]
Thread overview: 6+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2019-11-26 14:57 [PATCH, OpenACC, v3] Non-contiguous array support for OpenACC data clauses Chung-Lin Tang
2020-04-07 13:50 ` (v4 update) " Chung-Lin Tang
2023-03-10 13:24 ` Thomas Schwinge
2023-03-15 14:47 ` Thomas Schwinge
2023-03-24 15:17 ` Thomas Schwinge
2023-04-03 14:39 ` [og12] OpenACC: Pass pre-allocated 'ptrblock' to 'goacc_noncontig_array_create_ptrblock' [PR76739] (was: [PATCH, OpenACC, v3] Non-contiguous array support for OpenACC data clauses) Thomas Schwinge
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).