diff options
author | Chung-Lin Tang <cltang@codesourcery.com> | 2020-04-19 05:10:43 -0700 |
---|---|---|
committer | Kwok Cheung Yeung <kcy@codesourcery.com> | 2022-06-21 14:11:06 +0100 |
commit | 15d0f61a7fecdc8fd12857c40879ea3730f6d99f (patch) | |
tree | eb33062c1c8241e39819cd70f2f58f70db91ff8a /gcc | |
parent | 18bdbd7e50129beac70a58d582619be1deb63f45 (diff) | |
download | gcc-15d0f61a7fecdc8fd12857c40879ea3730f6d99f.zip gcc-15d0f61a7fecdc8fd12857c40879ea3730f6d99f.tar.gz gcc-15d0f61a7fecdc8fd12857c40879ea3730f6d99f.tar.bz2 |
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 <cltang@codesourcery.com>
PR other/76739
gcc/c/
* 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.
(handle_omp_array_sections): Adjust handle_omp_array_sections_1 call,
handle non-contiguous case to create dynamic array map.
gcc/cp/
* 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.
(handle_omp_array_sections): Adjust handle_omp_array_sections_1 call,
handle non-contiguous case to create dynamic array map.
gcc/fortran/
* f95-lang.cc (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.cc (gimplify_scan_omp_clauses): Skip gimplification of
OMP_CLAUSE_SIZE of non-contiguous array maps (which is a TREE_LIST).
* 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.
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.
Diffstat (limited to 'gcc')
-rw-r--r-- | gcc/ChangeLog.omp | 24 | ||||
-rw-r--r-- | gcc/builtin-types.def | 3 | ||||
-rw-r--r-- | gcc/c/ChangeLog.omp | 9 | ||||
-rw-r--r-- | gcc/c/c-typeck.cc | 45 | ||||
-rw-r--r-- | gcc/cp/ChangeLog.omp | 9 | ||||
-rw-r--r-- | gcc/cp/semantics.cc | 45 | ||||
-rw-r--r-- | gcc/fortran/ChangeLog.omp | 6 | ||||
-rw-r--r-- | gcc/fortran/f95-lang.cc | 13 | ||||
-rw-r--r-- | gcc/fortran/types.def | 3 | ||||
-rw-r--r-- | gcc/gimplify.cc | 12 | ||||
-rw-r--r-- | gcc/omp-builtins.def | 2 | ||||
-rw-r--r-- | gcc/omp-expand.cc | 13 | ||||
-rw-r--r-- | gcc/omp-low.cc | 261 | ||||
-rw-r--r-- | gcc/testsuite/ChangeLog.omp | 5 | ||||
-rw-r--r-- | gcc/testsuite/c-c++-common/goacc/noncontig_array-1.c | 25 | ||||
-rw-r--r-- | gcc/tree-pretty-print.cc | 36 |
16 files changed, 487 insertions, 24 deletions
diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp index 236d795..eed10a5 100644 --- a/gcc/ChangeLog.omp +++ b/gcc/ChangeLog.omp @@ -1,3 +1,27 @@ +2020-04-19 Chung-Lin Tang <cltang@codesourcery.com> + + 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.cc (gimplify_scan_omp_clauses): Skip gimplification of + OMP_CLAUSE_SIZE of non-contiguous array maps (which is a TREE_LIST). + * 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. + 2022-06-20 Kwok Cheung Yeung <kcy@codesourcery.com> * Makefile.in (REVISION_s): Change default message. diff --git a/gcc/builtin-types.def b/gcc/builtin-types.def index 3a7cecd..eb21aef 100644 --- a/gcc/builtin-types.def +++ b/gcc/builtin-types.def @@ -840,6 +840,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 0000000..76ed322 --- /dev/null +++ b/gcc/c/ChangeLog.omp @@ -0,0 +1,9 @@ +2020-04-19 Chung-Lin Tang <cltang@codesourcery.com> + + 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. + (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.cc b/gcc/c/c-typeck.cc index e130196..4913a0a 100644 --- a/gcc/c/c-typeck.cc +++ b/gcc/c/c-typeck.cc @@ -13207,12 +13207,14 @@ c_finish_omp_cancellation_point (location_t loc, tree clauses) <= FIRST_NON_ONE we diagnose non-contiguous arrays if low bound isn't 0 or length isn't the array domain max + 1, for > FIRST_NON_ONE we can if MAYBE_ZERO_LEN is false. MAYBE_ZERO_LEN will be true in the above - case though, as some lengths could be zero. */ + case though, as some lengths could be zero. + NON_CONTIGUOUS will be true if this is an OpenACC non-contiguous array + section. */ static tree handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types, bool &maybe_zero_len, unsigned int &first_non_one, - enum c_omp_region_type ort) + bool &non_contiguous, enum c_omp_region_type ort) { tree ret, low_bound, length, type; if (TREE_CODE (t) != TREE_LIST) @@ -13330,7 +13332,8 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types, } ret = handle_omp_array_sections_1 (c, TREE_CHAIN (t), types, - maybe_zero_len, first_non_one, ort); + maybe_zero_len, first_non_one, + non_contiguous, ort); if (ret == error_mark_node || ret == NULL_TREE) return ret; @@ -13557,7 +13560,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_CHAIN (t)) == TREE_LIST) @@ -13570,10 +13575,15 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types, tree d_length = TREE_VALUE (d); if (d_length == NULL_TREE || !integer_onep (d_length)) { - 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; + } } } } @@ -13604,6 +13614,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 @@ -13614,7 +13625,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) @@ -13648,6 +13659,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; @@ -13671,6 +13683,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)) @@ -13767,6 +13786,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 0000000..624388c --- /dev/null +++ b/gcc/cp/ChangeLog.omp @@ -0,0 +1,9 @@ +2020-04-19 Chung-Lin Tang <cltang@codesourcery.com> + + 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. + (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.cc b/gcc/cp/semantics.cc index ab48f11..8fa23df 100644 --- a/gcc/cp/semantics.cc +++ b/gcc/cp/semantics.cc @@ -5046,12 +5046,14 @@ omp_privatize_field (tree t, bool shared) <= FIRST_NON_ONE we diagnose non-contiguous arrays if low bound isn't 0 or length isn't the array domain max + 1, for > FIRST_NON_ONE we can if MAYBE_ZERO_LEN is false. MAYBE_ZERO_LEN will be true in the above - case though, as some lengths could be zero. */ + case though, as some lengths could be zero. + NON_CONTIGUOUS will be true if this is an OpenACC non-contiguous array + section. */ static tree handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types, bool &maybe_zero_len, unsigned int &first_non_one, - enum c_omp_region_type ort) + bool &non_contiguous, enum c_omp_region_type ort) { tree ret, low_bound, length, type; if (TREE_CODE (t) != TREE_LIST) @@ -5149,7 +5151,8 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types, && TREE_CODE (TREE_CHAIN (t)) == FIELD_DECL) TREE_CHAIN (t) = omp_privatize_field (TREE_CHAIN (t), false); ret = handle_omp_array_sections_1 (c, TREE_CHAIN (t), types, - maybe_zero_len, first_non_one, ort); + maybe_zero_len, first_non_one, + non_contiguous, ort); if (ret == error_mark_node || ret == NULL_TREE) return ret; @@ -5389,7 +5392,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_AFFINITY && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND && TREE_CODE (TREE_CHAIN (t)) == TREE_LIST) @@ -5402,10 +5407,15 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types, tree d_length = TREE_VALUE (d); if (d_length == NULL_TREE || !integer_onep (d_length)) { - 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; + } } } } @@ -5448,6 +5458,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 @@ -5458,7 +5469,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) @@ -5493,6 +5504,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; @@ -5518,6 +5530,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)) @@ -5609,6 +5628,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 0000000..b742e1b --- /dev/null +++ b/gcc/fortran/ChangeLog.omp @@ -0,0 +1,6 @@ +2020-04-19 Chung-Lin Tang <cltang@codesourcery.com> + + 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. diff --git a/gcc/fortran/f95-lang.cc b/gcc/fortran/f95-lang.cc index 1a895a2..4640f73 100644 --- a/gcc/fortran/f95-lang.cc +++ b/gcc/fortran/f95-lang.cc @@ -643,6 +643,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, \ @@ -665,6 +667,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 @@ -1156,6 +1159,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] \ @@ -1197,6 +1209,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 cd79ad4..d290620 100644 --- a/gcc/fortran/types.def +++ b/gcc/fortran/types.def @@ -276,6 +276,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 2588824..8daa552 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -9570,8 +9570,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 cfa6483..0926f56 100644 --- a/gcc/omp-builtins.def +++ b/gcc/omp-builtins.def @@ -30,7 +30,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 ee70831..8bace75 100644 --- a/gcc/omp-expand.cc +++ b/gcc/omp-expand.cc @@ -10252,6 +10252,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.cc b/gcc/omp-low.cc index bf5779b..e714e2c 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -961,6 +961,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 * @@ -1626,6 +1743,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) @@ -3048,6 +3197,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 @@ -3056,7 +3249,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); @@ -3069,6 +3261,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); @@ -12702,6 +12902,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)); @@ -12776,8 +12985,15 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) && is_gimple_omp_oacc (ctx->stmt) && 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 @@ -12974,6 +13190,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)) { @@ -13120,6 +13338,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)); @@ -13192,6 +13433,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) @@ -13469,6 +13713,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 0000000..64bb0cb --- /dev/null +++ b/gcc/testsuite/ChangeLog.omp @@ -0,0 +1,5 @@ +2020-04-19 Chung-Lin Tang <cltang@codesourcery.com> + + 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 0000000..ea738f5 --- /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.cc b/gcc/tree-pretty-print.cc index 99af977..d52dfa0 100644 --- a/gcc/tree-pretty-print.cc +++ b/gcc/tree-pretty-print.cc @@ -943,6 +943,33 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION: pp_string (pp, "attach_zero_length_array_section"); 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 (); } @@ -953,8 +980,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: |