From d7904bb2b5d31443b36a16a819574d3f1664a9c9 Mon Sep 17 00:00:00 2001 From: Chung-Lin Tang Date: Sun, 19 Apr 2020 05:10:43 -0700 Subject: [PATCH] Merge non-contiguous array support patches. This version is based from v4, posted upstream here: https://gcc.gnu.org/pipermail/gcc-patches/2020-April/543437.html 2020-04-19 Chung-Lin Tang 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. include/ * 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. --- gcc/ChangeLog.omp | 23 ++ gcc/builtin-types.def | 3 + gcc/c/ChangeLog.omp | 9 + gcc/c/c-typeck.c | 45 ++- gcc/cp/ChangeLog.omp | 9 + gcc/cp/semantics.c | 45 ++- gcc/fortran/ChangeLog.omp | 6 + gcc/fortran/f95-lang.c | 13 + gcc/fortran/types.def | 3 + gcc/gimplify.c | 12 +- gcc/omp-builtins.def | 2 +- gcc/omp-expand.c | 13 + gcc/omp-low.c | 261 +++++++++++++++++- gcc/testsuite/ChangeLog.omp | 5 + .../c-c++-common/goacc/noncontig_array-1.c | 25 ++ gcc/tree-pretty-print.c | 36 ++- include/ChangeLog.omp | 12 + include/gomp-constants.h | 23 ++ libgomp/ChangeLog.omp | 32 +++ libgomp/libgomp.h | 4 + libgomp/libgomp_g.h | 2 +- libgomp/oacc-int.h | 51 ++++ libgomp/oacc-parallel.c | 204 +++++++++++++- libgomp/target.c | 190 ++++++++++++- .../noncontig_array-1.c | 103 +++++++ .../noncontig_array-2.c | 37 +++ .../noncontig_array-3.c | 45 +++ .../noncontig_array-4.c | 36 +++ .../noncontig_array-utils.h | 44 +++ 29 files changed, 1253 insertions(+), 40 deletions(-) create mode 100644 gcc/ChangeLog.omp create mode 100644 gcc/c/ChangeLog.omp create mode 100644 gcc/cp/ChangeLog.omp create mode 100644 gcc/fortran/ChangeLog.omp create mode 100644 gcc/testsuite/ChangeLog.omp create mode 100644 gcc/testsuite/c-c++-common/goacc/noncontig_array-1.c create mode 100644 include/ChangeLog.omp create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-1.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-2.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-3.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-4.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-utils.h diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp new file mode 100644 index 000000000000..fe70dd7d1bfe --- /dev/null +++ b/gcc/ChangeLog.omp @@ -0,0 +1,23 @@ +2020-04-19 Chung-Lin Tang + + PR other/76739 + + * 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. diff --git a/gcc/builtin-types.def b/gcc/builtin-types.def index d160826e1d4f..9df843fc91b0 100644 --- a/gcc/builtin-types.def +++ b/gcc/builtin-types.def @@ -839,6 +839,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/ChangeLog.omp b/gcc/c/ChangeLog.omp new file mode 100644 index 000000000000..523e82a4bdcf --- /dev/null +++ b/gcc/c/ChangeLog.omp @@ -0,0 +1,9 @@ +2020-04-19 Chung-Lin Tang + + PR other/76739 + + * 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. diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index 51a62c800f79..96f5f777cae1 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -13050,12 +13050,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 &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) @@ -13148,7 +13150,8 @@ handle_omp_array_sections_1 (tree c, tree t, vec &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; @@ -13371,14 +13374,21 @@ handle_omp_array_sections_1 (tree c, tree t, vec &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 @@ -13407,6 +13417,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 types; tree *tp = &OMP_CLAUSE_DECL (c); if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND @@ -13416,7 +13427,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) @@ -13449,6 +13460,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; @@ -13472,6 +13484,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)) @@ -13568,6 +13587,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/ChangeLog.omp b/gcc/cp/ChangeLog.omp new file mode 100644 index 000000000000..6b63998284d1 --- /dev/null +++ b/gcc/cp/ChangeLog.omp @@ -0,0 +1,9 @@ +2020-04-19 Chung-Lin Tang + + PR other/76739 + + * 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. diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index 4520181d4e59..aefe3f97eb07 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -4929,12 +4929,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 &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) @@ -5020,7 +5022,8 @@ handle_omp_array_sections_1 (tree c, tree t, vec &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; @@ -5256,14 +5259,21 @@ handle_omp_array_sections_1 (tree c, tree t, vec &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 @@ -5303,6 +5313,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 types; tree *tp = &OMP_CLAUSE_DECL (c); if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND @@ -5312,7 +5323,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) @@ -5346,6 +5357,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; @@ -5371,6 +5383,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)) @@ -5462,6 +5481,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/ChangeLog.omp b/gcc/fortran/ChangeLog.omp new file mode 100644 index 000000000000..38e2f39512d6 --- /dev/null +++ b/gcc/fortran/ChangeLog.omp @@ -0,0 +1,6 @@ +2020-04-19 Chung-Lin Tang + + PR other/76739 + + * f95-lang.c (DEF_FUNCTION_TYPE_VAR_5): New symbol. + * types.def (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_VAR): New type. diff --git a/gcc/fortran/f95-lang.c b/gcc/fortran/f95-lang.c index a346457b5e08..b6eb29e12666 100644 --- a/gcc/fortran/f95-lang.c +++ b/gcc/fortran/f95-lang.c @@ -636,6 +636,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, \ @@ -658,6 +660,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 @@ -1149,6 +1152,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] \ @@ -1190,6 +1202,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 8626ed0a4f8e..bd51a22d7a06 100644 --- a/gcc/fortran/types.def +++ b/gcc/fortran/types.def @@ -274,6 +274,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 b65106b14596..00f56d3fbdbb 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -9020,8 +9020,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 cfbf1e67b8e5..c31a04f7198b 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 dc797f95154f..53fd276fbeda 100644 --- a/gcc/omp-expand.c +++ b/gcc/omp-expand.c @@ -9854,6 +9854,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 7b122059c6e8..6d3602154b3e 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -954,6 +954,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 * @@ -1519,6 +1636,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) @@ -2916,6 +3065,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 @@ -2924,7 +3117,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); @@ -2937,6 +3129,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); @@ -11915,6 +12115,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_FORCE_DETACH: gcc_assert (is_gimple_omp_oacc (stmt)); @@ -11988,8 +12197,15 @@ 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); 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 (var, rcv_by_ref, ctx); if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER @@ -12178,6 +12394,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) vec_alloc (vkind, map_cnt); unsigned int map_idx = 0; + vec nca_descrs = vNULL; + for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c)) switch (OMP_CLAUSE_CODE (c)) { @@ -12268,6 +12486,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)); @@ -12340,6 +12581,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) @@ -12593,6 +12837,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/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp new file mode 100644 index 000000000000..64bb0cb2e5ce --- /dev/null +++ b/gcc/testsuite/ChangeLog.omp @@ -0,0 +1,5 @@ +2020-04-19 Chung-Lin Tang + + PR other/76739 + + * c-c++-common/goacc/noncontig_array-1.c: New test. 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..ea738f5b65b4 --- /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 0a575eb9dad7..3a2f01bfe603 100644 --- a/gcc/tree-pretty-print.c +++ b/gcc/tree-pretty-print.c @@ -894,6 +894,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 (); } @@ -904,8 +931,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/ChangeLog.omp b/include/ChangeLog.omp new file mode 100644 index 000000000000..9bd3fb60a78d --- /dev/null +++ b/include/ChangeLog.omp @@ -0,0 +1,12 @@ +2020-04-19 Chung-Lin Tang + + 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. diff --git a/include/gomp-constants.h b/include/gomp-constants.h index 6e163b025609..3a4334ba0794 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) @@ -151,6 +152,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. */ @@ -186,6 +207,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/ChangeLog.omp b/libgomp/ChangeLog.omp index 568ad9ccac6f..5234168e4b1e 100644 --- a/libgomp/ChangeLog.omp +++ b/libgomp/ChangeLog.omp @@ -1,3 +1,35 @@ +2020-04-19 Chung-Lin Tang + + PR other/76739 + + * 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. + 2018-11-28 Julian Brown Tobias Burnus diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h index c1a8ee264cb2..7a71df21ba0e 100644 --- a/libgomp/libgomp.h +++ b/libgomp/libgomp.h @@ -1228,6 +1228,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 3cbe0a4ca89e..63031501c86e 100644 --- a/libgomp/libgomp_g.h +++ b/libgomp/libgomp_g.h @@ -378,7 +378,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 2cfafe7929c7..8ab73519aa8e 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 4f48b052239b..55e16427e732 100644 --- a/libgomp/oacc-parallel.c +++ b/libgomp/oacc-parallel.c @@ -36,7 +36,7 @@ #include #include #include - +#include /* 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,172 @@ _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 *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; +} /* Handle the mapping pair that are presented when a deviceptr clause is used with Fortran. */ @@ -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,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; @@ -290,8 +466,9 @@ 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); #ifdef RC_CHECKING gomp_mutex_lock (&acc_dev->lock); @@ -381,7 +558,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); @@ -472,8 +649,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; @@ -486,9 +662,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; diff --git a/libgomp/target.c b/libgomp/target.c index 9b952f79b7ca..e8cb6ec2cdee 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -852,10 +852,11 @@ 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; bool has_always_ptrset = false; const int rshift = short_mapkind ? 8 : 3; @@ -863,8 +864,9 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, 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; tgt->prev = NULL; @@ -1012,6 +1014,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]; @@ -1144,6 +1168,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, false, + /* 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) @@ -1409,6 +1471,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)) @@ -1567,6 +1638,100 @@ 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, + 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->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, + 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) @@ -1608,13 +1773,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 * @@ -1625,7 +1803,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 000000000000..a70375c03f4e --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-1.c @@ -0,0 +1,103 @@ +/* { dg-do run } */ + +#include +#include + +#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 +#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 +#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 +#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..554bda77bbda --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/noncontig_array-utils.h @@ -0,0 +1,44 @@ +#include +#include +#include +#include + +/* 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; +} -- 2.47.2