[gcc/devel/omp/gcc-16] Non-contiguous array support patches [PR76739]
Sandra Loosemore via Gcc-cvs <[email protected]>
| Newsgroups | gmane.comp.gcc.cvs |
|---|---|
| Message-ID | <[email protected]> |
https://gcc.gnu.org/g:4ecff95f061b1e5d9478b3c5d7643146d7364ecb commit 4ecff95f061b1e5d9478b3c5d7643146d7364ecb Author: Chung-Lin Tang <[email protected]> Date: Thu Apr 2 20:53:38 2026 +0000 Non-contiguous array support patches [PR76739] This is based on OG14 commit b143c1c447945ce05903ff1360ead97774dfce4b, which was based from v4, posted upstream here: https://gcc.gnu.org/pipermail/gcc-patches/2020-April/543437.html It also incorporates a number of follow-up bug and bit-rot fixes, OG14 commits e11726d3467543de45448097dde27ba34bf04bfe 87ea4de1c4a360d5d62357491a41811213f4528c 151fc161d0ed640048444ca18f9325e3d2e03e99 628a000bdbf63252c2ede13ccab8e99a19769866 11263c048d39ab1d6a11067b18674bf8307bbbf5 8c1068bbe3e52529bede5466a43af8d98f38dac2 and OG15 commits 130451de0f3 27cfb2b1363 1f60b10dd70 gcc/c/ChangeLog PR other/76739 * c-typeck.cc (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. Reject non-DECL base-pointer cases as unsupported. (handle_omp_array_sections): Adjust handle_omp_array_sections_1 call, handle non-contiguous case to create dynamic array map. gcc/cp/ChangeLog PR other/76739 * semantics.cc (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. Reject non-DECL base-pointer cases as unsupported. (handle_omp_array_sections): Adjust handle_omp_array_sections_1 call, handle non-contiguous case to create dynamic array map. gcc/fortran/ChangeLog PR other/76739 * f95-lang.cc (DEF_FUNCTION_TYPE_VAR_5): New symbol. * types.def (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR): New type. gcc/ChangeLog PR other/76739 * builtin-types.def (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR): New type. * gimplify.cc (omp_group_base): Handle GOMP_MAP_NONCONTIG_ARRAY_*. (gimplify_scan_omp_clauses): Handle OMP_TARGET_UPDATE. (gimplify_adjust_omp_clauses): Skip gimplification of OMP_CLAUSE_SIZE of non-contiguous array maps (which is a TREE_LIST). * omp-builtins.def (BUILT_IN_GOACC_DATA_START): Adjust function type to new BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR. * omp-expand.cc (expand_omp_target): Add non-contiguous array descriptor pointers to variadic arguments. * omp-low.cc (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. * tree-pretty-print.cc (dump_omp_clause): Handle GOMP_MAP_NONCONTIG_ARRAY_*. gcc/testsuite/ChangeLog PR other/76739 * c-c++-common/goacc/data-clause-1.c (foo): Remove expected message. * c-c++-common/goacc/noncontig_array-1.c: New test. * g++.dg/goacc/data-clause-1.C (foo): Remove expected message. include/ChangeLog PR other/76739 * gomp-constants.h (GOMP_MAP_FLAG_SPECIAL_3): Define. (enum gomp_map_kind): Add GOMP_MAP_NONCONTIG_ARRAY, GOMP_MAP_NONCONTIG_ARRAY_TO, GOMP_MAP_NONCONTIG_ARRAY_FROM, GOMP_MAP_NONCONTIG_ARRAY_TOFROM, GOMP_MAP_NONCONTIG_ARRAY_FORCE_TO, GOMP_MAP_NONCONTIG_ARRAY_FORCE_FROM, GOMP_MAP_NONCONTIG_ARRAY_FORCE_TOFROM, GOMP_MAP_NONCONTIG_ARRAY_ALLOC, GOMP_MAP_NONCONTIG_ARRAY_FORCE_ALLOC, GOMP_MAP_NONCONTIG_ARRAY_FORCE_PRESENT. (GOMP_MAP_NONCONTIG_ARRAY_P): Define. libgomp/ChangeLog PR other/76739 * libgomp.h (goacc_map_vars): Add 'struct goacc_ncarray_info *' formal parameter. * libgomp_g.h (GOACC_data_start): Add variadic '...' to 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-mem.c (acc_map_data, goacc_enter_datum) (goacc_enter_data_internal): Adjust. * 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. (goacc_map_vars): Adjust. * 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. Co-Authored-By: Kwok Cheung Yeung <[email protected]> Co-Authored-By: Thomas Schwinge <[email protected]> Co-Authored-By: Paul-Antoine Arras <[email protected]> Co-Authored-By: Sandra Loosemore <[email protected]> Diff: --- gcc/builtin-types.def | 3 + gcc/c/c-typeck.cc | 50 +++- gcc/cp/semantics.cc | 46 +++- gcc/fortran/f95-lang.cc | 13 ++ gcc/fortran/types.def | 3 + gcc/gimplify.cc | 19 +- gcc/omp-builtins.def | 2 +- gcc/omp-expand.cc | 13 ++ gcc/omp-low.cc | 260 ++++++++++++++++++++- gcc/testsuite/c-c++-common/goacc/data-clause-1.c | 2 +- .../c-c++-common/goacc/noncontig_array-1.c | 26 +++ gcc/testsuite/g++.dg/goacc/data-clause-1.C | 2 +- gcc/tree-pretty-print.cc | 36 ++- include/gomp-constants.h | 22 ++ libgomp/libgomp.h | 5 +- libgomp/libgomp_g.h | 2 +- libgomp/oacc-int.h | 52 +++++ libgomp/oacc-mem.c | 8 +- libgomp/oacc-parallel.c | 204 +++++++++++++++- libgomp/target.c | 188 ++++++++++++++- .../libgomp.oacc-c-c++-common/noncontig_array-1.c | 103 ++++++++ .../libgomp.oacc-c-c++-common/noncontig_array-2.c | 37 +++ .../libgomp.oacc-c-c++-common/noncontig_array-3.c | 45 ++++ .../libgomp.oacc-c-c++-common/noncontig_array-4.c | 36 +++ .../noncontig_array-utils.h | 44 ++++ 25 files changed, 1185 insertions(+), 36 deletions(-) diff --git a/gcc/builtin-types.def b/gcc/builtin-types.def index b08dd516f4d6..d73b9a786f30 100644 --- a/gcc/builtin-types.def +++ b/gcc/builtin-types.def @@ -1076,6 +1076,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.cc b/gcc/c/c-typeck.cc index 40ef7ef94aa7..f246c6c524f3 100644 --- a/gcc/c/c-typeck.cc +++ b/gcc/c/c-typeck.cc @@ -15811,12 +15811,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; bool openacc = (ort & C_ORT_ACC) != 0; @@ -15892,7 +15894,8 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types, } ret = handle_omp_array_sections_1 (c, TREE_OPERAND (t, 0), 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; @@ -16119,7 +16122,9 @@ 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 could be non-contiguous. */ + array-section-subscript, the array section could be non-contiguous. + Note that OpenACC does accept these kinds of non-contiguous pointer + based arrays. */ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_AFFINITY && TREE_CODE (TREE_OPERAND (t, 0)) == OMP_ARRAY_SECTION) @@ -16133,6 +16138,24 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types, tree d_length = TREE_OPERAND (d, 2); if (d_length == NULL_TREE || !integer_onep (d_length)) { + if (openacc && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP) + { + while (TREE_CODE (d) == OMP_ARRAY_SECTION) + d = TREE_OPERAND (d, 0); + if (DECL_P (d)) + { + /* Note that OpenACC does accept these kinds of + non-contiguous pointer based arrays. */ + non_contiguous = true; + break; + } + error_at (OMP_CLAUSE_LOCATION (c), + "base-pointer expression in %qs clause not " + "supported for non-contiguous arrays", + omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + return error_mark_node; + } + error_at (OMP_CLAUSE_LOCATION (c), "array section is not contiguous in %qs clause", omp_clause_code_name[OMP_CLAUSE_CODE (c)]); @@ -16167,6 +16190,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 @@ -16175,7 +16199,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) @@ -16209,6 +16233,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; @@ -16232,6 +16257,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)) @@ -16328,6 +16360,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.cc b/gcc/cp/semantics.cc index cf1a9049adc4..bafb558b2594 100644 --- a/gcc/cp/semantics.cc +++ b/gcc/cp/semantics.cc @@ -6050,12 +6050,14 @@ public: <= 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; bool openacc = (ort & C_ORT_ACC) != 0; @@ -6119,7 +6121,8 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types, && TREE_CODE (TREE_OPERAND (t, 0)) == FIELD_DECL) TREE_OPERAND (t, 0) = omp_privatize_field (TREE_OPERAND (t, 0), false); ret = handle_omp_array_sections_1 (c, TREE_OPERAND (t, 0), types, - maybe_zero_len, first_non_one, ort); + maybe_zero_len, first_non_one, + non_contiguous, ort); if (ret == error_mark_node || ret == NULL_TREE) return ret; @@ -6372,6 +6375,24 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types, tree d_length = TREE_OPERAND (d, 2); if (d_length == NULL_TREE || !integer_onep (d_length)) { + if (openacc && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP) + { + while (TREE_CODE (d) == OMP_ARRAY_SECTION) + d = TREE_OPERAND (d, 0); + if (DECL_P (d)) + { + /* Note that OpenACC does accept these kinds of + non-contiguous pointer based arrays. */ + non_contiguous = true; + break; + } + error_at (OMP_CLAUSE_LOCATION (c), + "base-pointer expression in %qs clause not " + "supported for non-contiguous arrays", + omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + return error_mark_node; + } + error_at (OMP_CLAUSE_LOCATION (c), "array section is not contiguous in %qs clause", omp_clause_code_name[OMP_CLAUSE_CODE (c)]); @@ -6418,6 +6439,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 @@ -6426,7 +6448,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) @@ -6461,6 +6483,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; @@ -6488,6 +6511,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)) @@ -6579,6 +6609,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.cc b/gcc/fortran/f95-lang.cc index b796bcf142d7..5e4bc420bf46 100644 --- a/gcc/fortran/f95-lang.cc +++ b/gcc/fortran/f95-lang.cc @@ -678,6 +678,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, \ @@ -700,6 +702,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 @@ -1217,6 +1220,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] \ @@ -1258,6 +1270,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 6f98e3e48ffa..2ebb1e6a71c6 100644 --- a/gcc/fortran/types.def +++ b/gcc/fortran/types.def @@ -286,6 +286,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.cc b/gcc/gimplify.cc index cf10c73efbc7..105bbb7b759f 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -11402,6 +11402,14 @@ omp_group_base (omp_mapping_group *grp, unsigned int *chained, case GOMP_MAP_ALWAYS_PRESENT_FROM: case GOMP_MAP_ALWAYS_PRESENT_TO: case GOMP_MAP_ALWAYS_PRESENT_TOFROM: + case GOMP_MAP_NONCONTIG_ARRAY_ALLOC: + case GOMP_MAP_NONCONTIG_ARRAY_FROM: + case GOMP_MAP_NONCONTIG_ARRAY_TO: + case GOMP_MAP_NONCONTIG_ARRAY_TOFROM: + case GOMP_MAP_NONCONTIG_ARRAY_FORCE_ALLOC: + case GOMP_MAP_NONCONTIG_ARRAY_FORCE_FROM: + case GOMP_MAP_NONCONTIG_ARRAY_FORCE_TO: + case GOMP_MAP_NONCONTIG_ARRAY_FORCE_TOFROM: case GOMP_MAP_ALLOC: case GOMP_MAP_RELEASE: case GOMP_MAP_DELETE: @@ -14001,6 +14009,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, || code == OMP_TARGET_DATA || code == OMP_TARGET_ENTER_DATA || code == OMP_TARGET_EXIT_DATA + || code == OMP_TARGET_UPDATE || code == OACC_DATA || code == OACC_KERNELS || code == OACC_PARALLEL @@ -16219,7 +16228,15 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, gimplify_omp_ctxp = ctx->outer_context; gimple_seq *seq_p; seq_p = enter_omp_iterator_loop_context (c, loops_seq_p, pre_p); - if (gimplify_expr (&OMP_CLAUSE_SIZE (c), seq_p, NULL, + 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), seq_p, NULL, is_gimple_val, fb_rvalue) == GS_ERROR) { gimplify_omp_ctxp = ctx; diff --git a/gcc/omp-builtins.def b/gcc/omp-builtins.def index d7440a42761b..4d6718ed39de 100644 --- a/gcc/omp-builtins.def +++ b/gcc/omp-builtins.def @@ -31,7 +31,7 @@ along with GCC; see the file COPYING3. If not see doesn't source those. */ 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_DATA, "GOACC_enter_data", diff --git a/gcc/omp-expand.cc b/gcc/omp-expand.cc index ca77fb5914aa..625dbf423d5f 100644 --- a/gcc/omp-expand.cc +++ b/gcc/omp-expand.cc @@ -10677,6 +10677,19 @@ expand_omp_target (struct omp_region *region) gsi_insert_before (&gsi, g, GSI_SAME_STMT); } + /* 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.cc b/gcc/omp-low.cc index 55820cc85d75..4da441d5b9fc 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -988,6 +988,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 * @@ -1726,6 +1843,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, c); + install_var_local (array_decl, ctx); + break; + } + if (DECL_P (decl)) { if (DECL_SIZE (decl) @@ -3186,6 +3335,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 @@ -3194,7 +3387,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); @@ -3207,6 +3399,14 @@ scan_omp_target (gomp_target *stmt, omp_context *outer_ctx) TYPE_NAME (ctx->record_type) = name; TYPE_ARTIFICIAL (ctx->record_type) = 1; + /* 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); + if (offloaded) { create_omp_child_function (ctx, false); @@ -12991,6 +13191,15 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) case GOMP_MAP_FORCE_TOFROM: 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_FORCE_DETACH: gcc_assert (is_gimple_omp_oacc (stmt)); @@ -13075,8 +13284,16 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) && is_gimple_omp_oacc (ctx->stmt) && OMP_CLAUSE_MAP_IN_REDUCTION (c))) { + tree var_type = TREE_TYPE (var); x = build_receiver_ref (c, true, ctx); tree new_var = lookup_decl (var, ctx); + 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 (c, rcv_by_ref, ctx); if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER @@ -13418,6 +13635,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)) { @@ -13602,6 +13821,29 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) avar = lower_omp_map_iterator_expr (avar, c, stmt); 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)); @@ -13674,6 +13916,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) @@ -14113,6 +14358,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); + } + if (!deep_map_cnt) { DECL_INITIAL (TREE_VEC_ELT (t, 1)) diff --git a/gcc/testsuite/c-c++-common/goacc/data-clause-1.c b/gcc/testsuite/c-c++-common/goacc/data-clause-1.c index b9871ec2cd58..59c3ec7cbd79 100644 --- a/gcc/testsuite/c-c++-common/goacc/data-clause-1.c +++ b/gcc/testsuite/c-c++-common/goacc/data-clause-1.c @@ -98,7 +98,7 @@ foo (int g[3][10], int h[4][8], int i[2][10], int j[][9], bar (&j2[0][0]); #pragma acc parallel copy(q[1:2]) ; - #pragma acc parallel copy(q[3:5][ :10]) /* { dg-error "array section is not contiguous" } */ + #pragma acc parallel copy(q[3:5][ :10]) ; #pragma acc parallel copy(r[3: ][2:1][1: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 000000000000..fe7480ab2e7b --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/noncontig_array-1.c @@ -0,0 +1,26 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +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/testsuite/g++.dg/goacc/data-clause-1.C b/gcc/testsuite/g++.dg/goacc/data-clause-1.C index ead71855ba2d..5597b99ff637 100644 --- a/gcc/testsuite/g++.dg/goacc/data-clause-1.C +++ b/gcc/testsuite/g++.dg/goacc/data-clause-1.C @@ -99,7 +99,7 @@ foo (int g[3][10], int h[4][8], int i[2][10], int j[][9], bar (&j2[0][0]); #pragma acc parallel copy(q[1:2]) ; - #pragma acc parallel copy(q[3:5][ :10]) /* { dg-error "array section is not contiguous" } */ + #pragma acc parallel copy(q[3:5][ :10]) ; #pragma acc parallel copy(r[3: ][2:1][1:2]) ; diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc index e64c3a7f8414..1a4a2ccc8a25 100644 --- a/gcc/tree-pretty-print.cc +++ b/gcc/tree-pretty-print.cc @@ -1159,6 +1159,33 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) case GOMP_MAP_USES_ALLOCATORS: pp_string (pp, "uses_allocators"); 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 (); } @@ -1169,8 +1196,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 ba6c028c1d12..e690d955bdda 100644 --- a/include/gomp-constants.h +++ b/include/gomp-constants.h @@ -188,6 +188,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), /* Like GOMP_MAP_ATTACH, but allow attaching to zero-length array sections (i.e. set to NULL when array section is not mapped) Currently only used @@ -261,6 +281,8 @@ enum gomp_map_kind (((X) & GOMP_MAP_FLAG_PRESENT) == GOMP_MAP_FLAG_PRESENT \ || (X) == GOMP_MAP_FORCE_PRESENT) +#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 d4d0434d719d..8066870e8a6b 100644 --- a/libgomp/libgomp.h +++ b/libgomp/libgomp.h @@ -1548,10 +1548,13 @@ extern bool 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, struct gomp_offload_session *); extern void goacc_unmap_vars (struct target_mem_desc *, bool, diff --git a/libgomp/libgomp_g.h b/libgomp/libgomp_g.h index a3097d5b0936..7616547322dc 100644 --- a/libgomp/libgomp_g.h +++ b/libgomp/libgomp_g.h @@ -415,7 +415,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 e346ed50120b..a08515c138ae 100644 --- a/libgomp/oacc-int.h +++ b/libgomp/oacc-int.h @@ -165,6 +165,58 @@ 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 *, void *); + + #ifdef HAVE_ATTRIBUTE_VISIBILITY # pragma GCC visibility pop #endif diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index f9ceb094d156..ae3161d93acd 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -421,7 +421,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, NULL); + &kinds, NULL, true, GOMP_MAP_VARS_ENTER_DATA, NULL); assert (tgt); assert (tgt->list_count == 1); splay_tree_key n = tgt->list[0].key; @@ -586,7 +586,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, NULL); + kinds, NULL, true, GOMP_MAP_VARS_ENTER_DATA, NULL); assert (tgt); assert (tgt->list_count == 1); n = tgt->list[0].key; @@ -1227,7 +1227,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_ __attribute__((unused)) = 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, NULL); assert (tgt_ == NULL); gomp_mutex_lock (&acc_dev->lock); @@ -1278,7 +1278,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, NULL); assert (tgt); diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c index 3770af510f29..fd4ab245f977 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. */ @@ -46,6 +46,171 @@ _Static_assert (GOACC_FLAGS_UNMARSHAL (GOMP_DEVICE_HOST_FALLBACK) == GOACC_FLAG_HOST_FALLBACK, "legacy GOMP_DEVICE_HOST_FALLBACK broken"); +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 *ptrblock, + void *tgt_ptrblock_addr) +{ + struct goacc_ncarray_descr_type *descr = nca->descr; + void **tgt_data_rows = nca->tgt_data_rows; + 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); +} /* Handle the mapping pair that are presented when a deviceptr clause is used with Fortran. */ @@ -115,6 +280,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", @@ -250,13 +416,22 @@ 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)) { k.host_start = (uintptr_t) fn; @@ -297,8 +472,12 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (void *), : gomp_offload_session_new (acc_dev, gomp_alloca)); struct target_mem_desc *tgt - = goacc_map_vars (acc_dev, aq, mapnum, hostaddrs, NULL, sizes, kinds, true, - GOMP_MAP_VARS_TARGET, session); + = goacc_map_vars (acc_dev, aq, mapnum, hostaddrs, NULL, sizes, kinds, + nca_info, true, GOMP_MAP_VARS_TARGET, session); + if (aq == NULL) + free (nca_info); + else + acc_dev->openacc.async.queue_callback_func (aq, free, nca_info); if (profiling_p) { @@ -370,7 +549,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); @@ -462,16 +641,27 @@ 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 = goacc_map_vars (NULL, NULL, 0, NULL, NULL, NULL, NULL, true, 0, NULL); + tgt = goacc_map_vars (NULL, NULL, 0, NULL, NULL, NULL, NULL, NULL, + true, 0, 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 = goacc_map_vars (acc_dev, NULL, mapnum, hostaddrs, NULL, sizes, kinds, - true, 0, NULL); + nca_info, true, 0, NULL); + 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 5f855ee2fc99..8d1a786a44bb 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -1224,8 +1224,8 @@ 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, - htab_t *refcount_set, + void *kinds, struct goacc_ncarray_info *nca_info, + bool short_mapkind, htab_t *refcount_set, enum gomp_map_vars_kind pragma_kind, struct gomp_offload_session *session) { @@ -1233,6 +1233,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, assert (/* SESSION must be present iff doing target offload. */ !!session == target_p); 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; bool has_always_ptrset = false; bool openmp_p = (pragma_kind & GOMP_MAP_VARS_OPENACC) == 0; @@ -1246,8 +1247,9 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, iterators_p = gomp_merge_iterator_maps (&mapnum, &hostaddrs, &sizes, &kinds, &iterator_count); 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; tgt->prev = NULL; @@ -1440,6 +1442,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].is_attach = 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]; @@ -1575,6 +1599,45 @@ 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, false, false, + /* TODO: cbuf? */ NULL, + refcount_set); + } + 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) @@ -1937,6 +2000,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)) @@ -2171,6 +2243,105 @@ 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, false, false, + cbufp, refcount_set); + } + 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->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->is_attach = 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, false, + 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 = 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) + /* Free once the transfer has completed. */ + devicep->openacc.async.queue_callback_func (aq, free, ptrblock); + else + free (ptrblock); + } + } + } } if (target_p) @@ -2264,8 +2435,8 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum, struct target_mem_desc *tgt; tgt = gomp_map_vars_internal (devicep, NULL, mapnum, hostaddrs, devaddrs, - sizes, kinds, short_mapkind, refcount_set, - pragma_kind, session); + sizes, kinds, NULL, short_mapkind, + refcount_set, pragma_kind, session); if (local_refcount_set) htab_free (local_refcount_set); @@ -2276,12 +2447,13 @@ 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, struct gomp_offload_session *session) { return gomp_map_vars_internal (devicep, aq, mapnum, hostaddrs, devaddrs, - sizes, kinds, short_mapkind, NULL, + sizes, kinds, nca_info, short_mapkind, NULL, GOMP_MAP_VARS_OPENACC | pragma_kind, session); } 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 000000000000..a70375c03f4e --- /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 000000000000..b85c6371f259 --- /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 000000000000..99db207493e4 --- /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 000000000000..6cfaf98d37e3 --- /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 000000000000..6900d1ffa675 --- /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; +}