From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1709) id B40C73852744; Mon, 13 Jun 2022 13:15:59 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org B40C73852744 Content-Type: text/plain; charset="us-ascii" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit From: Chung-Lin Tang To: gcc-cvs@gcc.gnu.org Subject: [gcc/devel/omp/gcc-11] openmp: Implement uses_allocators clause X-Act-Checkin: gcc X-Git-Author: Chung-Lin Tang X-Git-Refname: refs/heads/devel/omp/gcc-11 X-Git-Oldrev: dfe61b1ec5f5c7953bebb11b0dfb7b32e2c39b02 X-Git-Newrev: 729c88aa2c98547bcdee26166c88b670c5d421c5 Message-Id: <20220613131559.B40C73852744@sourceware.org> Date: Mon, 13 Jun 2022 13:15:59 +0000 (GMT) X-BeenThere: gcc-cvs@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-cvs mailing list List-Unsubscribe: , List-Archive: List-Help: List-Subscribe: , X-List-Received-Date: Mon, 13 Jun 2022 13:15:59 -0000 https://gcc.gnu.org/g:729c88aa2c98547bcdee26166c88b670c5d421c5 commit 729c88aa2c98547bcdee26166c88b670c5d421c5 Author: Chung-Lin Tang Date: Mon Jun 13 21:12:32 2022 +0800 openmp: Implement uses_allocators clause This is a merge of: https://gcc.gnu.org/pipermail/gcc-patches/2022-June/596412.html For user defined allocator handles, this allows target regions to assign memory space and traits to allocators, and automatically calls omp_init/destroy_allocator() in the beginning/end of the target region. For pre-defined allocators (i.e. omp_..._mem_alloc names), this is a no-op, such clauses are not created. Asides from the front-end portions, the target region transforms are done in gimplify_omp_workshare. This patch also includes added changes to enforce the "allocate allocator must be also in a uses_allocator clause". This is done during gimplify_scan_omp_clauses. gcc/c-family/ChangeLog: * c-omp.cc (c_omp_split_clauses): Add OMP_CLAUSE_USES_ALLOCATORS case. * c-pragma.h (enum pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_USES_ALLOCATORS. gcc/c/ChangeLog: * c-parser.cc (c_parser_omp_clause_name): Add case for uses_allocators clause. (c_parser_omp_clause_uses_allocators): New function. (c_parser_omp_all_clauses): Add PRAGMA_OMP_CLAUSE_USES_ALLOCATORS case. (OMP_TARGET_CLAUSE_MASK): Add PRAGMA_OMP_CLAUSE_USES_ALLOCATORS to mask. * c-typeck.cc (c_finish_omp_clauses): Add case handling for OMP_CLAUSE_USES_ALLOCATORS. gcc/cp/ChangeLog: * parser.cc (cp_parser_omp_clause_name): Add case for uses_allocators clause. (cp_parser_omp_clause_uses_allocators): New function. (cp_parser_omp_all_clauses): Add PRAGMA_OMP_CLAUSE_USES_ALLOCATORS case. (OMP_TARGET_CLAUSE_MASK): Add PRAGMA_OMP_CLAUSE_USES_ALLOCATORS to mask. * semantics.cc (finish_omp_clauses): Add case handling for OMP_CLAUSE_USES_ALLOCATORS. fortran/ChangeLog: * gfortran.h (struct gfc_omp_namelist): Add memspace_sym, traits_sym fields. (OMP_LIST_USES_ALLOCATORS): New list enum. * openmp.cc (enum omp_mask2): Add OMP_CLAUSE_USES_ALLOCATORS. (gfc_match_omp_clause_uses_allocators): New function. (gfc_match_omp_clauses): Add case to handle OMP_CLAUSE_USES_ALLOCATORS. (OMP_TARGET_CLAUSES): Add OMP_CLAUSE_USES_ALLOCATORS. (resolve_omp_clauses): Add "USES_ALLOCATORS" to clause_names[]. * dump-parse-tree.cc (show_omp_namelist): Handle OMP_LIST_USES_ALLOCATORS. (show_omp_clauses): Likewise. * trans-array.cc (gfc_conv_array_initializer): Adjust array index to always be a created tree expression instead of NULL_TREE when zero. * trans-openmp.cc (gfc_trans_omp_clauses): For ALLOCATE clause, handle using gfc_trans_omp_variable for EXPR_VARIABLE exprs. Add handling of OMP_LIST_USES_ALLOCATORS case. * types.def (BT_FN_VOID_PTRMODE): Define. (BT_FN_PTRMODE_PTRMODE_INT_PTR): Define. gcc/ChangeLog: * builtin-types.def (BT_FN_VOID_PTRMODE): Define. (BT_FN_PTRMODE_PTRMODE_INT_PTR): Define. * omp-builtins.def (BUILT_IN_OMP_INIT_ALLOCATOR): Define. (BUILT_IN_OMP_DESTROY_ALLOCATOR): Define. * tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_USES_ALLOCATORS. * tree-pretty-print.cc (dump_omp_clause): Handle OMP_CLAUSE_USES_ALLOCATORS. * tree.h (OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR): New macro. (OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE): New macro. (OMP_CLAUSE_USES_ALLOCATORS_TRAITS): New macro. * tree.cc (omp_clause_num_ops): Add OMP_CLAUSE_USES_ALLOCATORS. (omp_clause_code_name): Add "uses_allocators". * gimplify.cc (gimplify_scan_omp_clauses): Add checking of OpenMP target region allocate clauses, to require a uses_allocators clause to exist for allocators. (gimplify_omp_workshare): Add handling of OMP_CLAUSE_USES_ALLOCATORS for OpenMP target regions; create calls of omp_init/destroy_allocator around target region body. * omp-low.cc (lower_private_allocate): Adjust receiving of allocator. (lower_rec_input_clauses): Likewise. (create_task_copyfn): Add dereference for allocator if needed. * system.h (startswith): New function. gcc/testsuite/ChangeLog: * c-c++-common/gomp/uses_allocators-1.c: New test. * c-c++-common/gomp/uses_allocators-2.c: New test. * c-c++-common/gomp/uses_allocators-3.c: New test. * gfortran.dg/gomp/allocate-1.f90: Adjust testcase. * gfortran.dg/gomp/uses_allocators-1.f90: New test. * gfortran.dg/gomp/uses_allocators-2.f90: New test. * gfortran.dg/gomp/uses_allocators-3.f90: New test. Diff: --- gcc/builtin-types.def | 3 + gcc/c-family/c-omp.c | 1 + gcc/c-family/c-pragma.h | 1 + gcc/c/c-parser.c | 216 ++++++++- gcc/c/c-typeck.c | 113 +++++ gcc/cp/parser.c | 237 ++++++++- gcc/cp/semantics.c | 107 +++++ gcc/fortran/dump-parse-tree.c | 15 + gcc/fortran/gfortran.h | 3 + gcc/fortran/openmp.c | 531 ++++++++++++++++++++- gcc/fortran/trans-array.c | 9 +- gcc/fortran/trans-openmp.c | 38 +- gcc/fortran/types.def | 3 + gcc/gimplify.c | 111 +++++ gcc/omp-builtins.def | 4 + gcc/omp-low.c | 13 +- gcc/system.h | 8 + .../c-c++-common/gomp/uses_allocators-1.c | 46 ++ .../c-c++-common/gomp/uses_allocators-2.c | 39 ++ .../c-c++-common/gomp/uses_allocators-3.c | 31 ++ gcc/testsuite/gfortran.dg/gomp/allocate-1.f90 | 3 +- .../gfortran.dg/gomp/uses_allocators-1.f90 | 53 ++ .../gfortran.dg/gomp/uses_allocators-2.f90 | 52 ++ .../gfortran.dg/gomp/uses_allocators-3.f90 | 14 + gcc/tree-core.h | 5 +- gcc/tree-pretty-print.c | 14 + gcc/tree.c | 3 + gcc/tree.h | 9 + 28 files changed, 1666 insertions(+), 16 deletions(-) diff --git a/gcc/builtin-types.def b/gcc/builtin-types.def index c4654cd7bfb..1094cb98b6e 100644 --- a/gcc/builtin-types.def +++ b/gcc/builtin-types.def @@ -283,6 +283,7 @@ DEF_FUNCTION_TYPE_1 (BT_FN_DFLOAT32_DFLOAT32, BT_DFLOAT32, BT_DFLOAT32) DEF_FUNCTION_TYPE_1 (BT_FN_DFLOAT64_DFLOAT64, BT_DFLOAT64, BT_DFLOAT64) DEF_FUNCTION_TYPE_1 (BT_FN_DFLOAT128_DFLOAT128, BT_DFLOAT128, BT_DFLOAT128) DEF_FUNCTION_TYPE_1 (BT_FN_VOID_VPTR, BT_VOID, BT_VOLATILE_PTR) +DEF_FUNCTION_TYPE_1 (BT_FN_VOID_PTRMODE, BT_VOID, BT_PTRMODE) DEF_FUNCTION_TYPE_1 (BT_FN_VOID_PTRPTR, BT_VOID, BT_PTR_PTR) DEF_FUNCTION_TYPE_1 (BT_FN_VOID_CONST_PTR, BT_VOID, BT_CONST_PTR) DEF_FUNCTION_TYPE_1 (BT_FN_UINT_UINT, BT_UINT, BT_UINT) @@ -642,6 +643,8 @@ DEF_FUNCTION_TYPE_3 (BT_FN_PTR_SIZE_SIZE_PTRMODE, BT_PTR, BT_SIZE, BT_SIZE, BT_PTRMODE) DEF_FUNCTION_TYPE_3 (BT_FN_VOID_PTR_UINT8_PTRMODE, BT_VOID, BT_PTR, BT_UINT8, BT_PTRMODE) +DEF_FUNCTION_TYPE_3 (BT_FN_PTRMODE_PTRMODE_INT_PTR, BT_PTRMODE, BT_PTRMODE, + BT_INT, BT_PTR) DEF_FUNCTION_TYPE_4 (BT_FN_SIZE_CONST_PTR_SIZE_SIZE_FILEPTR, BT_SIZE, BT_CONST_PTR, BT_SIZE, BT_SIZE, BT_FILEPTR) diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c index fed909e1da3..a426be960a3 100644 --- a/gcc/c-family/c-omp.c +++ b/gcc/c-family/c-omp.c @@ -2009,6 +2009,7 @@ c_omp_split_clauses (location_t loc, enum tree_code code, case OMP_CLAUSE_IS_DEVICE_PTR: case OMP_CLAUSE_DEFAULTMAP: case OMP_CLAUSE_DEPEND: + case OMP_CLAUSE_USES_ALLOCATORS: s = C_OMP_CLAUSE_SPLIT_TARGET; break; case OMP_CLAUSE_NUM_TEAMS: diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h index 145260e0c20..4ef13d104be 100644 --- a/gcc/c-family/c-pragma.h +++ b/gcc/c-family/c-pragma.h @@ -154,6 +154,7 @@ enum pragma_omp_clause { PRAGMA_OMP_CLAUSE_UNTIED, PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR, PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR, + PRAGMA_OMP_CLAUSE_USES_ALLOCATORS, /* Clauses for OpenACC. */ PRAGMA_OACC_CLAUSE_ASYNC, diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index 9a3d0cb8cea..47e78935fb6 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -12863,6 +12863,8 @@ c_parser_omp_clause_name (c_parser *parser) result = PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR; else if (!strcmp ("use_device_ptr", p)) result = PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR; + else if (!strcmp ("uses_allocators", p)) + result = PRAGMA_OMP_CLAUSE_USES_ALLOCATORS; break; case 'v': if (!strcmp ("vector", p)) @@ -15569,6 +15571,213 @@ c_parser_omp_clause_allocate (c_parser *parser, tree list) return nl; } +/* OpenMP 5.0: + uses_allocators ( allocator-list ) + + allocator-list: + allocator + allocator , allocator-list + allocator ( traits-array ) + allocator ( traits-array ) , allocator-list + + OpenMP 5.2: + + uses_allocators ( modifier : allocator ) + uses_allocators ( modifier , modifier : allocator ) + + modifier: + traits ( traits-array ) + memspace ( mem-space-handle ) */ + +static tree +c_parser_omp_clause_uses_allocators (c_parser *parser, tree list) +{ + location_t clause_loc = c_parser_peek_token (parser)->location; + tree t = NULL_TREE, nl = list; + matching_parens parens; + if (!parens.require_open (parser)) + return list; + + tree memspace_expr = NULL_TREE; + tree traits_var = NULL_TREE; + + struct item_tok + { + location_t loc; + tree id; + item_tok (void) : loc (UNKNOWN_LOCATION), id (NULL_TREE) {} + }; + struct item { item_tok name, arg; }; + auto_vec *modifiers = NULL, *allocators = NULL; + auto_vec *cur_list = new auto_vec (4); + + while (true) + { + item it; + + if (c_parser_next_token_is (parser, CPP_NAME)) + { + c_token *tok = c_parser_peek_token (parser); + it.name.id = tok->value; + it.name.loc = tok->location; + c_parser_consume_token (parser); + + if (c_parser_next_token_is (parser, CPP_OPEN_PAREN)) + { + matching_parens parens2; + parens2.consume_open (parser); + + if (c_parser_next_token_is (parser, CPP_NAME)) + { + tok = c_parser_peek_token (parser); + it.arg.id = tok->value; + it.arg.loc = tok->location; + c_parser_consume_token (parser); + } + else + { + c_parser_error (parser, "expected identifier"); + parens2.skip_until_found_close (parser); + goto end; + } + parens2.skip_until_found_close (parser); + } + } + + cur_list->safe_push (it); + + if (c_parser_next_token_is (parser, CPP_COMMA)) + c_parser_consume_token (parser); + else if (c_parser_next_token_is (parser, CPP_COLON)) + { + if (modifiers) + { + c_parser_error (parser, "expected %<)%>"); + goto end; + } + else + { + c_parser_consume_token (parser); + modifiers = cur_list; + cur_list = new auto_vec (4); + } + } + else if (c_parser_next_token_is (parser, CPP_CLOSE_PAREN)) + { + gcc_assert (allocators == NULL); + allocators = cur_list; + cur_list = NULL; + break; + } + else + { + c_parser_error (parser, "expected %<)%>"); + goto end; + } + } + + if (modifiers) + for (unsigned i = 0; i < modifiers->length (); i++) + { + item& it = (*modifiers)[i]; + const char *p = IDENTIFIER_POINTER (it.name.id); + int strcmp_traits = 1, strcmp_memspace = 1; + + if ((strcmp_traits = strcmp ("traits", p)) == 0 + || (strcmp_memspace = strcmp ("memspace", p)) == 0) + { + if ((strcmp_traits == 0 && traits_var != NULL_TREE) + || (strcmp_memspace == 0 && memspace_expr != NULL_TREE)) + { + error_at (it.name.loc, "duplicate %qs modifier", p); + goto end; + } + t = lookup_name (it.arg.id); + if (t == NULL_TREE) + { + undeclared_variable (it.arg.loc, it.arg.id); + t = error_mark_node; + } + else if (strcmp_memspace == 0) + memspace_expr = t; + else if (strcmp_traits == 0) + traits_var = t; + else + gcc_unreachable (); + } + else + { + error_at (it.name.loc, "unknown modifier %qE", it.name.id); + goto end; + } + } + + if (allocators) + { + if (modifiers) + { + if (allocators->length () > 1) + { + error_at ((*allocators)[1].name.loc, + "% clause only accepts a single " + "allocator when using modifiers"); + goto end; + } + else if ((*allocators)[0].arg.id) + { + error_at ((*allocators)[0].arg.loc, + "legacy %<%E(%E)%> traits syntax not allowed in " + "% clause when using modifiers", + (*allocators)[0].name.id, (*allocators)[0].arg.id); + goto end; + } + } + + for (unsigned i = 0; i < allocators->length (); i++) + { + item& it = (*allocators)[i]; + t = lookup_name (it.name.id); + if (t == NULL_TREE) + { + undeclared_variable (it.name.loc, it.name.id); + goto end; + } + else if (t != error_mark_node) + { + tree t2 = NULL_TREE; + if (it.arg.id) + { + t2 = lookup_name (it.arg.id); + if (t2 == NULL_TREE) + { + undeclared_variable (it.arg.loc, it.arg.id); + goto end; + } + } + else + t2 = traits_var; + + tree c = build_omp_clause (clause_loc, + OMP_CLAUSE_USES_ALLOCATORS); + OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c) = t; + OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c) = memspace_expr; + OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c) = t2; + OMP_CLAUSE_CHAIN (c) = nl; + nl = c; + } + } + } + end: + if (cur_list) + delete cur_list; + if (modifiers) + delete modifiers; + if (allocators) + delete allocators; + parens.skip_until_found_close (parser); + return nl; +} + /* OpenMP 4.0: linear ( variable-list ) linear ( variable-list : expression ) @@ -16997,6 +17206,10 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask, clauses = c_parser_omp_clause_allocate (parser, clauses); c_name = "allocate"; break; + case PRAGMA_OMP_CLAUSE_USES_ALLOCATORS: + clauses = c_parser_omp_clause_uses_allocators (parser, clauses); + c_name = "uses_allocators"; + break; case PRAGMA_OMP_CLAUSE_LINEAR: clauses = c_parser_omp_clause_linear (parser, clauses); c_name = "linear"; @@ -21021,7 +21234,8 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser, | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEFAULTMAP) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IN_REDUCTION) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_THREAD_LIMIT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR)) + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR)\ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USES_ALLOCATORS)) static bool c_parser_omp_target (c_parser *parser, enum pragma_context context, bool *if_p) diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index 7b40f5ae8df..d9df42ca492 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -14710,6 +14710,119 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) } break; + case OMP_CLAUSE_USES_ALLOCATORS: + t = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c); + if ((VAR_P (t) || TREE_CODE (t) == PARM_DECL) + && (bitmap_bit_p (&generic_head, DECL_UID (t)) + || bitmap_bit_p (&map_head, DECL_UID (t)) + || bitmap_bit_p (&firstprivate_head, DECL_UID (t)) + || bitmap_bit_p (&lastprivate_head, DECL_UID (t)))) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qE appears more than once in data clauses", t); + remove = true; + break; + } + else + bitmap_set_bit (&generic_head, DECL_UID (t)); + if (TREE_CODE (TREE_TYPE (t)) != ENUMERAL_TYPE + || strcmp (IDENTIFIER_POINTER (TYPE_IDENTIFIER (TREE_TYPE (t))), + "omp_allocator_handle_t") != 0) + { + error_at (OMP_CLAUSE_LOCATION (c), + "allocator must be of % type"); + remove = true; + break; + } + if (TREE_CODE (t) == CONST_DECL) + { + /* Currently for pre-defined allocators in libgomp, we do not + require additional init/fini inside target regions, so discard + such clauses. */ + remove = true; + + if (strcmp (IDENTIFIER_POINTER (DECL_NAME (t)), + "omp_null_allocator") == 0) + { + error_at (OMP_CLAUSE_LOCATION (c), + "% cannot be used in " + "% clause"); + break; + } + + if (OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c) + || OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c)) + { + error_at (OMP_CLAUSE_LOCATION (c), + "modifiers cannot be used with pre-defined " + "allocators"); + break; + } + } + t = OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c); + if (t != NULL_TREE + && (TREE_CODE (t) != CONST_DECL + || TREE_CODE (TREE_TYPE (t)) != ENUMERAL_TYPE + || strcmp (IDENTIFIER_POINTER (TYPE_IDENTIFIER (TREE_TYPE (t))), + "omp_memspace_handle_t") != 0)) + { + error_at (OMP_CLAUSE_LOCATION (c), "memspace modifier must be " + "constant enum of % type"); + remove = true; + break; + } + t = OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c); + if (t != NULL_TREE) + { + bool type_err = false; + + if (TREE_CODE (TREE_TYPE (t)) != ARRAY_TYPE + || DECL_SIZE (t) == NULL_TREE) + type_err = true; + else + { + tree elem_t = TREE_TYPE (TREE_TYPE (t)); + if (TREE_CODE (elem_t) != RECORD_TYPE + || strcmp (IDENTIFIER_POINTER (TYPE_IDENTIFIER (elem_t)), + "omp_alloctrait_t") != 0 + || !TYPE_READONLY (elem_t)) + type_err = true; + } + if (type_err) + { + error_at (OMP_CLAUSE_LOCATION (c), "traits array must be of " + "% type"); + remove = true; + } + else + { + tree cst_val = decl_constant_value_1 (t, true); + if (cst_val == t) + { + error_at (OMP_CLAUSE_LOCATION (c), "traits array must be " + "of constant values"); + + remove = true; + } + } + } + + if (remove) + break; + else + { + /* Create a private clause for the allocator variable, placed + prior to current uses_allocators clause. */ + tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (c), + OMP_CLAUSE_PRIVATE); + OMP_CLAUSE_DECL (nc) = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c); + OMP_CLAUSE_CHAIN (nc) = c; + *pc = nc; + + pc = &OMP_CLAUSE_CHAIN (c); + continue; + } + case OMP_CLAUSE_DEPEND: t = OMP_CLAUSE_DECL (c); if (t == NULL_TREE) diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index 3a9ea272f10..8bf97a86164 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -36003,6 +36003,8 @@ cp_parser_omp_clause_name (cp_parser *parser) result = PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR; else if (!strcmp ("use_device_ptr", p)) result = PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR; + else if (!strcmp ("uses_allocators", p)) + result = PRAGMA_OMP_CLAUSE_USES_ALLOCATORS; break; case 'v': if (!strcmp ("vector", p)) @@ -38249,6 +38251,234 @@ cp_parser_omp_clause_allocate (cp_parser *parser, tree list) return nlist; } +/* OpenMP 5.0: + uses_allocators ( allocator-list ) + + allocator-list: + allocator + allocator , allocator-list + allocator ( traits-array ) + allocator ( traits-array ) , allocator-list + + OpenMP 5.2: + + uses_allocators ( modifier : allocator ) + uses_allocators ( modifier , modifier : allocator ) + + modifier: + traits ( traits-array ) + memspace ( mem-space-handle ) */ + +static tree +cp_parser_omp_clause_uses_allocators (cp_parser *parser, tree list) +{ + location_t clause_loc + = cp_lexer_peek_token (parser->lexer)->location; + tree t = NULL_TREE, nl = list; + matching_parens parens; + if (!parens.require_open (parser)) + return list; + + tree memspace_expr = NULL_TREE; + tree traits_var = NULL_TREE; + + struct item_tok + { + location_t loc; + tree id; + item_tok (void) : loc (UNKNOWN_LOCATION), id (NULL_TREE) {} + }; + struct item { item_tok name, arg; }; + auto_vec *modifiers = NULL, *allocators = NULL; + auto_vec *cur_list = new auto_vec (4); + + while (true) + { + item it; + + if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)) + { + cp_token *tok = cp_lexer_peek_token (parser->lexer); + it.name.id = tok->u.value; + it.name.loc = tok->location; + cp_lexer_consume_token (parser->lexer); + + if (cp_lexer_next_token_is (parser->lexer, CPP_OPEN_PAREN)) + { + matching_parens parens2; + parens2.consume_open (parser); + + if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)) + { + tok = cp_lexer_peek_token (parser->lexer); + it.arg.id = tok->u.value; + it.arg.loc = tok->location; + cp_lexer_consume_token (parser->lexer); + } + else + { + cp_parser_error (parser, "expected identifier"); + cp_parser_skip_to_closing_parenthesis (parser, + /*recovering=*/true, + /*or_comma=*/false, + /*consume_paren=*/true); + goto end; + } + cp_parser_skip_to_closing_parenthesis (parser, + /*recovering=*/false, + /*or_comma=*/false, + /*consume_paren=*/true); + } + } + + cur_list->safe_push (it); + + if (cp_lexer_next_token_is (parser->lexer, CPP_COMMA)) + cp_lexer_consume_token (parser->lexer); + else if (cp_lexer_next_token_is (parser->lexer, CPP_COLON)) + { + if (modifiers) + { + cp_parser_error (parser, "expected %<)%>"); + cp_parser_skip_to_closing_parenthesis (parser, + /*recovering=*/true, + /*or_comma=*/false, + /*consume_paren=*/true); + goto end; + } + else + { + cp_lexer_consume_token (parser->lexer); + modifiers = cur_list; + cur_list = new auto_vec (4); + } + } + else if (cp_lexer_next_token_is (parser->lexer, CPP_CLOSE_PAREN)) + { + gcc_assert (allocators == NULL); + allocators = cur_list; + cur_list = NULL; + break; + } + else + { + cp_parser_error (parser, "expected %<)%>"); + cp_parser_skip_to_closing_parenthesis (parser, + /*recovering=*/true, + /*or_comma=*/false, + /*consume_paren=*/true); + goto end; + } + } + + if (modifiers) + for (unsigned i = 0; i < modifiers->length (); i++) + { + item& it = (*modifiers)[i]; + const char *p = IDENTIFIER_POINTER (it.name.id); + int strcmp_traits = 1, strcmp_memspace = 1; + + if ((strcmp_traits = strcmp ("traits", p)) == 0 + || (strcmp_memspace = strcmp ("memspace", p)) == 0) + { + if ((strcmp_traits == 0 && traits_var != NULL_TREE) + || (strcmp_memspace == 0 && memspace_expr != NULL_TREE)) + { + error_at (it.name.loc, "duplicate %qs modifier", p); + goto end; + } + t = cp_parser_lookup_name_simple (parser, it.arg.id, it.arg.loc); + if (t == error_mark_node) + { + cp_parser_name_lookup_error (parser, it.arg.id, t, NLE_NULL, + it.arg.loc); + } + else if (strcmp_memspace == 0) + memspace_expr = t; + else if (strcmp_traits == 0) + traits_var = t; + else + gcc_unreachable (); + } + else + { + error_at (it.name.loc, "unknown modifier %qE", it.name.id); + goto end; + } + } + + if (allocators) + { + if (modifiers) + { + if (allocators->length () > 1) + { + error_at ((*allocators)[1].name.loc, + "% clause only accepts a single " + "allocator when using modifiers"); + goto end; + } + else if ((*allocators)[0].arg.id) + { + error_at ((*allocators)[0].arg.loc, + "legacy %<%E(%E)%> traits syntax not allowed in " + "% clause when using modifiers", + (*allocators)[0].name.id, (*allocators)[0].arg.id); + goto end; + } + } + + for (unsigned i = 0; i < allocators->length (); i++) + { + item& it = (*allocators)[i]; + t = cp_parser_lookup_name_simple (parser, it.name.id, it.name.loc); + if (t == error_mark_node) + { + cp_parser_name_lookup_error (parser, it.name.id, t, NLE_NULL, + it.name.loc); + goto end; + } + else if (t != error_mark_node) + { + tree t2 = NULL_TREE; + if (it.arg.id) + { + t2 = cp_parser_lookup_name_simple (parser, it.arg.id, + it.arg.loc); + if (t2 == error_mark_node) + { + cp_parser_name_lookup_error (parser, it.arg.id, t2, + NLE_NULL, it.arg.loc); + goto end; + } + } + else + t2 = traits_var; + + tree c = build_omp_clause (clause_loc, + OMP_CLAUSE_USES_ALLOCATORS); + OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c) = t; + OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c) = memspace_expr; + OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c) = t2; + OMP_CLAUSE_CHAIN (c) = nl; + nl = c; + } + } + } + end: + if (cur_list) + delete cur_list; + if (modifiers) + delete modifiers; + if (allocators) + delete allocators; + cp_parser_skip_to_closing_parenthesis (parser, + /*recovering=*/false, + /*or_comma=*/false, + /*consume_paren=*/true); + return nl; +} + /* OpenMP 2.5: lastprivate ( variable-list ) @@ -39794,6 +40024,10 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses = cp_parser_omp_clause_allocate (parser, clauses); c_name = "allocate"; break; + case PRAGMA_OMP_CLAUSE_USES_ALLOCATORS: + clauses = cp_parser_omp_clause_uses_allocators (parser, clauses); + c_name = "uses_allocators"; + break; case PRAGMA_OMP_CLAUSE_LINEAR: { bool declare_simd = false; @@ -43811,7 +44045,8 @@ cp_parser_omp_target_update (cp_parser *parser, cp_token *pragma_tok, | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_ALLOCATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IN_REDUCTION) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_THREAD_LIMIT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR)) + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR)\ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USES_ALLOCATORS)) static bool cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index a066873b7f0..929738285d9 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -7731,6 +7731,113 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) } goto handle_field_decl; + case OMP_CLAUSE_USES_ALLOCATORS: + t = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c); + if (TREE_CODE (t) == FIELD_DECL) + { + sorry_at (OMP_CLAUSE_LOCATION (c), "class members not yet " + "supported in % clause"); + remove = true; + break; + } + t = convert_from_reference (t); + if (TREE_CODE (TREE_TYPE (t)) != ENUMERAL_TYPE + || strcmp (IDENTIFIER_POINTER (TYPE_IDENTIFIER (TREE_TYPE (t))), + "omp_allocator_handle_t") != 0) + { + error_at (OMP_CLAUSE_LOCATION (c), + "allocator must be of % type"); + remove = true; + break; + } + if (TREE_CODE (t) == CONST_DECL) + { + /* Currently for pre-defined allocators in libgomp, we do not + require additional init/fini inside target regions, so discard + such clauses. */ + remove = true; + + if (strcmp (IDENTIFIER_POINTER (DECL_NAME (t)), + "omp_null_allocator") == 0) + { + error_at (OMP_CLAUSE_LOCATION (c), + "% cannot be used in " + "% clause"); + break; + } + + if (OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c) + || OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c)) + { + error_at (OMP_CLAUSE_LOCATION (c), + "modifiers cannot be used with pre-defined " + "allocators"); + break; + } + } + t = OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c); + if (t != NULL_TREE + && (TREE_CODE (t) != CONST_DECL + || TREE_CODE (TREE_TYPE (t)) != ENUMERAL_TYPE + || strcmp (IDENTIFIER_POINTER (TYPE_IDENTIFIER (TREE_TYPE (t))), + "omp_memspace_handle_t") != 0)) + { + error_at (OMP_CLAUSE_LOCATION (c), "memspace modifier must be " + "constant enum of % type"); + remove = true; + break; + } + t = OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c); + if (t != NULL_TREE) + { + bool type_err = false; + + if (TREE_CODE (TREE_TYPE (t)) != ARRAY_TYPE + || DECL_SIZE (t) == NULL_TREE) + type_err = true; + else + { + tree elem_t = TREE_TYPE (TREE_TYPE (t)); + if (TREE_CODE (elem_t) != RECORD_TYPE + || strcmp (IDENTIFIER_POINTER (TYPE_IDENTIFIER (elem_t)), + "omp_alloctrait_t") != 0 + || !TYPE_READONLY (elem_t)) + type_err = true; + } + if (type_err) + { + error_at (OMP_CLAUSE_LOCATION (c), "traits array %qE must be of " + "% type", t); + remove = true; + } + else + { + tree cst_val = decl_constant_value (t); + if (cst_val == t) + { + error_at (OMP_CLAUSE_LOCATION (c), "traits array must be " + "of constant values"); + + remove = true; + } + } + } + if (remove) + break; + else + { + /* Create a private clause for the allocator variable, placed + prior to current uses_allocators clause. */ + tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (c), + OMP_CLAUSE_PRIVATE); + OMP_CLAUSE_DECL (nc) = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c); + OMP_CLAUSE_CHAIN (nc) = c; + *pc = nc; + + pc = &OMP_CLAUSE_CHAIN (c); + continue; + } + case OMP_CLAUSE_DEPEND: t = OMP_CLAUSE_DECL (c); if (t == NULL_TREE) diff --git a/gcc/fortran/dump-parse-tree.c b/gcc/fortran/dump-parse-tree.c index a53accfd54e..65995ad53cd 100644 --- a/gcc/fortran/dump-parse-tree.c +++ b/gcc/fortran/dump-parse-tree.c @@ -1421,6 +1421,20 @@ show_omp_namelist (int list_type, gfc_omp_namelist *n) case OMP_LINEAR_UVAL: fputs ("uval(", dumpfile); break; default: break; } + else if (list_type == OMP_LIST_USES_ALLOCATORS) + { + show_symbol (n->sym); + fputs ("(memspace:", dumpfile); + if (n->memspace_sym) + show_symbol (n->traits_sym); + fputs (",traits:", dumpfile); + if (n->memspace_sym) + show_symbol (n->traits_sym); + fputc (')', dumpfile); + if (n->next) + fputc (',', dumpfile); + continue; + } fprintf (dumpfile, "%s", n->sym->name); if (list_type == OMP_LIST_LINEAR && n->u.linear_op != OMP_LINEAR_DEFAULT) fputc (')', dumpfile); @@ -1686,6 +1700,7 @@ show_omp_clauses (gfc_omp_clauses *omp_clauses) case OMP_LIST_ALLOCATE: type = "ALLOCATE"; break; case OMP_LIST_SCAN_IN: type = "INCLUSIVE"; break; case OMP_LIST_SCAN_EX: type = "EXCLUSIVE"; break; + case OMP_LIST_USES_ALLOCATORS: type = "USES_ALLOCATORS"; break; default: gcc_unreachable (); } diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h index 2de4507189c..13fef16754c 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -1352,6 +1352,8 @@ typedef struct gfc_omp_namelist struct gfc_omp_namelist_udr *udr; gfc_namespace *ns; } u2; + struct gfc_symbol *memspace_sym; + struct gfc_symbol *traits_sym; struct gfc_omp_namelist *next; locus where; } @@ -1393,6 +1395,7 @@ enum OMP_LIST_NONTEMPORAL, OMP_LIST_ALLOCATE, OMP_LIST_ALLOCATOR, + OMP_LIST_USES_ALLOCATORS, OMP_LIST_NUM }; diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c index ece04c03a68..e08d94ee060 100644 --- a/gcc/fortran/openmp.c +++ b/gcc/fortran/openmp.c @@ -980,6 +980,7 @@ enum omp_mask2 OMP_CLAUSE_ATTACH, OMP_CLAUSE_NOHOST, OMP_CLAUSE_ALLOCATOR, + OMP_CLAUSE_USES_ALLOCATORS, /* This must come last. */ OMP_MASK2_LAST }; @@ -1396,6 +1397,528 @@ gfc_match_omp_clause_reduction (char pc, gfc_omp_clauses *c, bool openacc, return MATCH_YES; } +/* OpenMP 5.0: + uses_allocators ( allocator-list ) + + allocator-list: + allocator + allocator , allocator-list + allocator ( traits-array ) + allocator ( traits-array ) , allocator-list + + OpenMP 5.2: + + uses_allocators ( modifier : allocator ) + uses_allocators ( modifier , modifier : allocator ) + + modifier: + traits ( traits-array ) + memspace ( mem-space-handle ) */ + +static match +gfc_match_omp_clause_uses_allocators (gfc_omp_clauses *c) +{ + char buffer[GFC_MAX_SYMBOL_LEN + 1]; + gfc_symbol *sym; + gfc_symbol *memspace_sym= NULL; + gfc_symbol *traits_sym= NULL; + locus traits_sym_loc; + match m, ret = MATCH_ERROR; + + if (gfc_match ("uses_allocators ( ") != MATCH_YES) + return MATCH_NO; + + struct item_tok + { + locus loc; + char *str; + item_tok (void) : str (NULL) {} + ~item_tok (void) { if (str) free (str); } + }; + struct item { item_tok name, arg; }; + auto_vec *modifiers = NULL, *allocators = NULL; + auto_vec *cur_list = new auto_vec (4); + + gfc_symbol *allocator_handle_kind; + + if (gfc_find_symbol ("omp_allocator_handle_kind", NULL, 1, &sym) + || sym == NULL + || sym->attr.dimension + || sym->value == NULL + || sym->value->expr_type != EXPR_CONSTANT + || sym->value->ts.type != BT_INTEGER) + { + gfc_error ("OpenMP % constant not found by " + "% clause at %C"); + goto error; + } + allocator_handle_kind = sym; + + while (true) + { + item it; + + m = gfc_match_name (buffer); + if (m == MATCH_YES) + { + it.name.str = xstrdup (buffer); + it.name.loc = gfc_current_locus; + } + else + { + gfc_error ("Expected identifier at %C"); + goto error; + } + + if (gfc_match_char ('(') == MATCH_YES) + { + m = gfc_match_name (buffer); + if (m == MATCH_YES) + { + it.arg.str = xstrdup (buffer); + it.arg.loc = gfc_current_locus; + } + else + { + gfc_error ("Expected identifier at %C"); + goto error; + } + if (gfc_match_char (')') != MATCH_YES) + { + gfc_error ("Expected %<)%> at %C"); + goto error; + } + } + + cur_list->safe_push (it); + it.name.str = NULL; + it.arg.str = NULL; + + if (gfc_match (" , ") == MATCH_YES) + continue; + else if (gfc_match (" : ") == MATCH_YES) + { + if (modifiers) + { + gfc_error ("expected %<)%> at %C"); + goto error; + } + else + { + modifiers = cur_list; + cur_list = new auto_vec (4); + } + } + else if (gfc_match_char (')') == MATCH_YES) + { + gcc_assert (allocators == NULL); + allocators = cur_list; + cur_list = NULL; + break; + } + else + { + gfc_error ("expected %<)%> at %C"); + goto error; + } + } + + if (modifiers) + for (unsigned i = 0; i < modifiers->length (); i++) + { + item& it = (*modifiers)[i]; + const char *p = it.name.str; + int strcmp_traits = 1, strcmp_memspace = 1; + gfc_symbol *sym; + + if ((strcmp_traits = strcmp ("traits", p)) == 0 + || (strcmp_memspace = strcmp ("memspace", p)) == 0) + { + if ((strcmp_traits == 0 && traits_sym != NULL) + || (strcmp_memspace == 0 && memspace_sym != NULL)) + { + gfc_error ("duplicate %qs modifier at %L", p, &it.name.loc); + goto error; + } + if (gfc_find_symbol (it.arg.str, NULL, 1, &sym) || sym == NULL) + { + gfc_error ("Symbol %qs at %L is ambiguous", + it.arg.str, &it.arg.loc); + goto error; + } + else if (strcmp_memspace == 0) + { + memspace_sym = sym; + + /* We have a memspace specified, now check if it is valid. + Start with finding if we have the standards specified + 'omp_memspace_handle_kind' available. */ + if (gfc_find_symbol ("omp_memspace_handle_kind", NULL, 1, &sym) + || sym == NULL + || sym->attr.dimension + || sym->value == NULL + || sym->value->expr_type != EXPR_CONSTANT + || sym->value->ts.type != BT_INTEGER) + { + gfc_error ("OpenMP % constant " + "not found by % clause at %L", + &it.arg.loc); + goto error; + } + + gfc_symbol *memspace_handle_kind = sym; + + if (memspace_sym->ts.type != BT_INTEGER + || memspace_sym->attr.flavor != FL_PARAMETER + || mpz_cmp_si (memspace_handle_kind->value->value.integer, + memspace_sym->ts.kind) != 0 + /* Check if identifier is of 'omp_..._mem_space' format. */ + || !startswith (memspace_sym->name, "omp_") + || !endswith (memspace_sym->name, "_mem_space")) + { + gfc_error ("%<%s%> at %L is not a pre-defined memory space " + "name", memspace_sym->name, &it.arg.loc); + goto error; + } + } + else if (strcmp_traits == 0) + { + traits_sym = sym; + traits_sym_loc = it.arg.loc; + } + else + gcc_unreachable (); + } + else + { + gfc_error ("unknown modifier %qs at %L", p, &it.name.loc); + goto error; + } + } + + if (allocators) + { + if (modifiers) + { + if (allocators->length () > 1) + { + gfc_error ("% clause only accepts a single " + "allocator when using modifiers at %L", + &(*allocators)[1].name.loc); + goto error; + } + else if ((*allocators)[0].arg.str) + { + gfc_error ("legacy %<%s(%s)%> traits syntax not allowed in " + "% clause when using modifiers at %L", + (*allocators)[0].name.str, (*allocators)[0].arg.str, + &(*allocators)[0].arg.loc); + goto error; + } + } + + for (unsigned i = 0; i < allocators->length (); i++) + { + item& it = (*allocators)[i]; + + gfc_symbol *allocator_sym; + locus allocator_sym_loc; + + if (gfc_find_symbol (it.name.str, NULL, 1, &allocator_sym) != 0 + || allocator_sym == NULL) + { + gfc_error ("Symbol %qs at %L is ambiguous", + it.name.str, &it.name.loc); + goto error; + } + allocator_sym_loc = it.name.loc; + + gfc_symbol *curr_traits_sym; + locus curr_traits_sym_loc; + + if (it.arg.str) + { + if (gfc_find_symbol (it.arg.str, NULL, 1, &curr_traits_sym) + || curr_traits_sym == NULL) + { + gfc_error ("Symbol %qs at %L is ambiguous", + it.arg.str, &it.arg.loc); + goto error; + } + curr_traits_sym_loc = it.arg.loc; + } + else + { + curr_traits_sym = traits_sym; + curr_traits_sym_loc = traits_sym_loc; + } + + if (curr_traits_sym) + { + if (curr_traits_sym->ts.type != BT_DERIVED + || strcmp (curr_traits_sym->ts.u.derived->name, + "omp_alloctrait") != 0 + || curr_traits_sym->attr.flavor != FL_PARAMETER + || curr_traits_sym->as->rank != 1) + { + gfc_error ("%<%s%> at %L must be of constant " + "% array type and have a " + "constant initializer", curr_traits_sym->name, + &curr_traits_sym_loc); + goto error; + } + gfc_set_sym_referenced (curr_traits_sym); + } + + if (allocator_sym->ts.type != BT_INTEGER + || mpz_cmp_si (allocator_handle_kind->value->value.integer, + allocator_sym->ts.kind) != 0) + { + gfc_error ("%<%s%> at %L must be integer of %<%s%> kind", + allocator_sym->name, &allocator_sym_loc, + allocator_handle_kind->name); + goto error; + } + + if (allocator_sym->attr.flavor == FL_PARAMETER) + { + if (strcmp (allocator_sym->name, "omp_null_allocator") == 0) + { + gfc_error ("% cannot be used in " + "% clause at %L", + &allocator_sym_loc); + goto error; + } + + /* Check if identifier is a 'omp_..._mem_alloc' pre-defined + allocator. */ + if (!startswith (allocator_sym->name, "omp_") + || !endswith (allocator_sym->name, "_mem_alloc")) + { + gfc_error ("%<%s%> at %L is not a pre-defined memory " + "allocator", allocator_sym->name, + &allocator_sym_loc); + goto error; + } + + /* Currently for pre-defined allocators in libgomp, we do not + require additional init/fini inside target regions, + so do nothing here to discard such clauses. */ + } + else + { + gfc_set_sym_referenced (allocator_sym); + + gfc_omp_namelist *n = gfc_get_omp_namelist (); + n->sym = allocator_sym; + n->memspace_sym = memspace_sym; + n->traits_sym = curr_traits_sym; + n->where = it.name.loc; + + n->next = c->lists[OMP_LIST_USES_ALLOCATORS]; + c->lists[OMP_LIST_USES_ALLOCATORS] = n; + } + } + } + + ret = MATCH_YES; + + end: + if (cur_list) + delete cur_list; + if (modifiers) + delete modifiers; + if (allocators) + delete allocators; + return ret; + + error: + ret = MATCH_ERROR; + gfc_error_check (); + goto end; + +#if 0 + do + { + if (++i > 2) + { + gfc_error ("Only two modifiers are allowed on % " + "clause at %C"); + goto error; + } + + if (gfc_match ("memspace ( ") == MATCH_YES) + { + if (memspace_seen) + { + gfc_error ("Multiple memspace modifiers at %C"); + goto error; + } + memspace_seen = true; + m = gfc_match_symbol (&sym, 1); + if (m == MATCH_YES) + memspace_sym = sym; + else + goto error; + if (gfc_match_char (')') != MATCH_YES) + goto error; + } + else if (gfc_match ("traits ( ") == MATCH_YES) + { + if (traits_seen) + { + gfc_error ("Multiple traits modifiers at %C"); + goto error; + } + traits_seen = true; + m = gfc_match_symbol (&sym, 1); + if (m == MATCH_YES) + traits_sym = sym; + else + goto error; + if (gfc_match_char (')') != MATCH_YES) + goto error; + } + else + break; + } + while (gfc_match (" , ") == MATCH_YES); + + if ((memspace_seen || traits_seen) + && gfc_match (" : ") != MATCH_YES) + goto error; + + while (true) + { + m = gfc_match_symbol (&sym, 1); + if (m != MATCH_YES) + { + gfc_error ("Expected name of allocator at %C"); + goto error; + } + gfc_symbol *allocator_sym = sym; + + if (gfc_match_char ('(') == MATCH_YES) + { + if (memspace_seen || traits_seen) + { + gfc_error ("Modifiers cannot be used with (deprecated) traits " + "array list syntax at %C"); + goto error; + } + m = gfc_match_symbol (&sym, 1); + if (m == MATCH_YES) + traits_sym = sym; + else + goto error; + if (gfc_match_char (')') != MATCH_YES) + goto error; + } + + if (traits_sym) + { + if (traits_sym->ts.type != BT_DERIVED + || strcmp (traits_sym->ts.u.derived->name, + "omp_alloctrait") != 0 + || traits_sym->attr.flavor != FL_PARAMETER + || traits_sym->as->rank != 1) + { + gfc_error ("%<%s%> at %C must be of constant " + "% array type and have a " + "constant initializer", traits_sym->name); + goto error; + } + gfc_set_sym_referenced (traits_sym); + } + + if (memspace_sym) + { + if (gfc_find_symbol ("omp_memspace_handle_kind", NULL, 1, &sym) + || sym == NULL + || sym->attr.dimension + || sym->value == NULL + || sym->value->expr_type != EXPR_CONSTANT + || sym->value->ts.type != BT_INTEGER) + { + gfc_error ("OpenMP % constant not " + "found by % clause at %C"); + goto error; + } + gfc_symbol *memspace_handle_kind = sym; + + if (memspace_sym->ts.type != BT_INTEGER + || memspace_sym->attr.flavor != FL_PARAMETER + || mpz_cmp_si (memspace_handle_kind->value->value.integer, + memspace_sym->ts.kind) != 0 + /* Check if identifier is of 'omp_..._mem_space' format. */ + || !startswith (memspace_sym->name, "omp_") + || !endswith (memspace_sym->name, "_mem_space")) + { + gfc_error ("%<%s%> at %C is not a pre-defined memory space name", + memspace_sym->name); + goto error; + } + } + + if (allocator_sym->ts.type != BT_INTEGER + || mpz_cmp_si (allocator_handle_kind->value->value.integer, + allocator_sym->ts.kind) != 0) + { + gfc_error ("%<%s%> at %C must be integer of %<%s%> kind", + allocator_sym->name, allocator_handle_kind->name); + goto error; + } + + if (allocator_sym->attr.flavor == FL_PARAMETER) + { + /* Check if identifier is a 'omp_..._mem_alloc' pre-defined + allocator. */ + if (!startswith (allocator_sym->name, "omp_") + || !endswith (allocator_sym->name, "_mem_alloc")) + { + gfc_error ("%<%s%> at %C is not a pre-defined memory allocator", + allocator_sym->name); + goto error; + } + + /* Currently for pre-defined allocators in libgomp, we do not + require additional init/fini inside target regions, + so do nothing here to discard such clauses. */ + } + else + { + gfc_set_sym_referenced (allocator_sym); + + gfc_omp_namelist *n = gfc_get_omp_namelist (); + n->sym = allocator_sym; + n->memspace_sym = memspace_sym; + n->traits_sym = traits_sym; + n->where = gfc_current_locus; + + n->next = c->lists[OMP_LIST_USES_ALLOCATORS]; + c->lists[OMP_LIST_USES_ALLOCATORS] = n; + } + + if (gfc_match (" , ") == MATCH_YES) + { + if (memspace_seen || traits_seen) + { + gfc_error ("When using modifiers, only a single allocator can be " + "specified in each % clause at %C"); + goto error; + } + } + else + break; + + memspace_sym = NULL; + traits_sym = NULL; + } + + if (gfc_match_char (')') != MATCH_YES) + goto error; +#endif +} /* Match with duplicate check. Matches 'name'. If expr != NULL, it then matches '(expr)', otherwise, if open_parens is true, @@ -2950,6 +3473,9 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, ("use_device_addr (", &c->lists[OMP_LIST_USE_DEVICE_ADDR], false) == MATCH_YES) continue; + if ((mask & OMP_CLAUSE_USES_ALLOCATORS) + && gfc_match_omp_clause_uses_allocators (c) == MATCH_YES) + continue; break; case 'v': /* VECTOR_LENGTH must be matched before VECTOR, because the latter @@ -3661,7 +4187,8 @@ cleanup: | OMP_CLAUSE_DEPEND | OMP_CLAUSE_NOWAIT | OMP_CLAUSE_PRIVATE \ | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULTMAP \ | OMP_CLAUSE_IS_DEVICE_PTR | OMP_CLAUSE_IN_REDUCTION \ - | OMP_CLAUSE_THREAD_LIMIT | OMP_CLAUSE_ALLOCATE) + | OMP_CLAUSE_THREAD_LIMIT | OMP_CLAUSE_ALLOCATE \ + | OMP_CLAUSE_USES_ALLOCATORS) #define OMP_TARGET_DATA_CLAUSES \ (omp_mask (OMP_CLAUSE_DEVICE) | OMP_CLAUSE_MAP | OMP_CLAUSE_IF \ | OMP_CLAUSE_USE_DEVICE_PTR | OMP_CLAUSE_USE_DEVICE_ADDR) @@ -6529,7 +7056,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, "IN_REDUCTION", "TASK_REDUCTION", "DEVICE_RESIDENT", "LINK", "USE_DEVICE", "CACHE", "IS_DEVICE_PTR", "USE_DEVICE_PTR", "USE_DEVICE_ADDR", - "NONTEMPORAL", "ALLOCATE", "ALLOCATOR" }; + "NONTEMPORAL", "ALLOCATE", "ALLOCATOR", "USES_ALLOCATORS" }; STATIC_ASSERT (ARRAY_SIZE (clause_names) == OMP_LIST_NUM); if (omp_clauses == NULL) diff --git a/gcc/fortran/trans-array.c b/gcc/fortran/trans-array.c index 8e5277594ab..41b28cb6843 100644 --- a/gcc/fortran/trans-array.c +++ b/gcc/fortran/trans-array.c @@ -6462,10 +6462,6 @@ gfc_conv_array_initializer (tree type, gfc_expr * expr) &expr->where, flag_max_array_constructor); return NULL_TREE; } - if (mpz_cmp_si (c->offset, 0) != 0) - index = gfc_conv_mpz_to_tree (c->offset, gfc_index_integer_kind); - else - index = NULL_TREE; if (mpz_cmp_si (c->repeat, 1) > 0) { @@ -6490,6 +6486,11 @@ gfc_conv_array_initializer (tree type, gfc_expr * expr) else range = NULL; + if (range == NULL || mpz_cmp_si (c->offset, 0) != 0) + index = gfc_conv_mpz_to_tree (c->offset, gfc_index_integer_kind); + else + index = NULL_TREE; + gfc_init_se (&se, NULL); switch (c->expr->expr_type) { diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index a51d8227c76..c3d662e5c46 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -4069,9 +4069,16 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, if (n->expr) { tree allocator_; - gfc_init_se (&se, NULL); - gfc_conv_expr (&se, n->expr); - allocator_ = gfc_evaluate_now (se.expr, block); + if (n->expr->expr_type == EXPR_VARIABLE) + allocator_ + = gfc_trans_omp_variable (n->expr->symtree->n.sym, + false); + else + { + gfc_init_se (&se, NULL); + gfc_conv_expr (&se, n->expr); + allocator_ = gfc_evaluate_now (se.expr, block); + } OMP_CLAUSE_ALLOCATE_ALLOCATOR (node) = allocator_; } omp_clauses = gfc_trans_add_clause (node, omp_clauses); @@ -5153,6 +5160,29 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, omp_clauses = gfc_trans_add_clause (node, omp_clauses); } break; + case OMP_LIST_USES_ALLOCATORS: + for (; n != NULL; n = n->next) + { + tree allocator = gfc_trans_omp_variable (n->sym, false); + tree memspace = (n->memspace_sym + ? gfc_conv_constant_to_tree (n->memspace_sym->value) + : NULL_TREE); + tree traits = (n->traits_sym + ? gfc_trans_omp_variable (n->traits_sym, false) + : NULL_TREE); + + tree nc = build_omp_clause (input_location, OMP_CLAUSE_PRIVATE); + OMP_CLAUSE_DECL (nc) = allocator; + omp_clauses = gfc_trans_add_clause (nc, omp_clauses); + + nc = build_omp_clause (input_location, + OMP_CLAUSE_USES_ALLOCATORS); + OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (nc) = allocator; + OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (nc) = memspace; + OMP_CLAUSE_USES_ALLOCATORS_TRAITS (nc) = traits; + omp_clauses = gfc_trans_add_clause (nc, omp_clauses); + } + break; default: break; } @@ -7606,6 +7636,8 @@ gfc_split_omp_clauses (gfc_code *code, = code->ext.omp_clauses->device; clausesa[GFC_OMP_SPLIT_TARGET].thread_limit = code->ext.omp_clauses->thread_limit; + clausesa[GFC_OMP_SPLIT_TARGET].lists[OMP_LIST_USES_ALLOCATORS] + = code->ext.omp_clauses->lists[OMP_LIST_USES_ALLOCATORS]; for (int i = 0; i < OMP_DEFAULTMAP_CAT_NUM; i++) clausesa[GFC_OMP_SPLIT_TARGET].defaultmap[i] = code->ext.omp_clauses->defaultmap[i]; diff --git a/gcc/fortran/types.def b/gcc/fortran/types.def index bb0c15c9fe9..2ce5cabac0d 100644 --- a/gcc/fortran/types.def +++ b/gcc/fortran/types.def @@ -81,6 +81,7 @@ DEF_FUNCTION_TYPE_0 (BT_FN_UINT, BT_UINT) DEF_FUNCTION_TYPE_0 (BT_FN_VOID, BT_VOID) DEF_FUNCTION_TYPE_1 (BT_FN_VOID_PTR, BT_VOID, BT_PTR) +DEF_FUNCTION_TYPE_1 (BT_FN_VOID_PTRMODE, BT_VOID, BT_PTRMODE) DEF_FUNCTION_TYPE_1 (BT_FN_VOID_PTRPTR, BT_VOID, BT_PTR_PTR) DEF_FUNCTION_TYPE_1 (BT_FN_VOID_VPTR, BT_VOID, BT_VOLATILE_PTR) DEF_FUNCTION_TYPE_1 (BT_FN_INT_INT, BT_INT, BT_INT) @@ -155,6 +156,8 @@ DEF_FUNCTION_TYPE_3 (BT_FN_VOID_SIZE_SIZE_PTR, BT_VOID, BT_SIZE, BT_SIZE, DEF_FUNCTION_TYPE_3 (BT_FN_UINT_UINT_PTR_PTR, BT_UINT, BT_UINT, BT_PTR, BT_PTR) DEF_FUNCTION_TYPE_3 (BT_FN_PTR_SIZE_SIZE_PTRMODE, BT_PTR, BT_SIZE, BT_SIZE, BT_PTRMODE) +DEF_FUNCTION_TYPE_3 (BT_FN_PTRMODE_PTRMODE_INT_PTR, BT_PTRMODE, BT_PTRMODE, + BT_INT, BT_PTR) DEF_FUNCTION_TYPE_4 (BT_FN_VOID_OMPFN_PTR_UINT_UINT, BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR, BT_UINT, BT_UINT) diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 6fc6a375b11..da71ed4d5d0 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -10955,6 +10955,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, case OMP_CLAUSE_BIND: case OMP_CLAUSE_IF_PRESENT: case OMP_CLAUSE_FINALIZE: + case OMP_CLAUSE_USES_ALLOCATORS: break; case OMP_CLAUSE_ORDER: @@ -11069,6 +11070,49 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, remove = true; break; } + if ((omp_requires_mask & OMP_REQUIRES_DYNAMIC_ALLOCATORS) == 0 + && OMP_CLAUSE_ALLOCATE_ALLOCATOR (c) + && TREE_CODE (OMP_CLAUSE_ALLOCATE_ALLOCATOR (c)) != INTEGER_CST) + { + tree allocator = OMP_CLAUSE_ALLOCATE_ALLOCATOR (c); + tree clauses = NULL_TREE; + + /* Get clause list of the nearest enclosing target construct. */ + if (ctx->code == OMP_TARGET) + clauses = *orig_list_p; + else + { + struct gimplify_omp_ctx *tctx = ctx->outer_context; + while (tctx && tctx->code != OMP_TARGET) + tctx = tctx->outer_context; + if (tctx) + clauses = tctx->clauses; + } + + if (clauses) + { + tree uc; + if (TREE_CODE (allocator) == MEM_REF + || TREE_CODE (allocator) == INDIRECT_REF) + allocator = TREE_OPERAND (allocator, 0); + for (uc = clauses; uc; uc = OMP_CLAUSE_CHAIN (uc)) + if (OMP_CLAUSE_CODE (uc) == OMP_CLAUSE_USES_ALLOCATORS) + { + tree uc_allocator + = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (uc); + if (operand_equal_p (allocator, uc_allocator)) + break; + } + if (uc == NULL_TREE) + { + error_at (OMP_CLAUSE_LOCATION (c), "allocator %<%qE%> " + "requires % clause in " + "target region", allocator, allocator); + remove = true; + break; + } + } + } if (gimplify_expr (&OMP_CLAUSE_ALLOCATE_ALLOCATOR (c), pre_p, NULL, is_gimple_val, fb_rvalue) == GS_ERROR) { @@ -14559,6 +14603,73 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p) body = NULL; gimple_seq_add_stmt (&body, g); } + else if ((ort & ORT_TARGET) != 0 && (ort & ORT_ACC) == 0) + { + gimple_seq init_seq = NULL; + gimple_seq fini_seq = NULL; + + tree omp_init_allocator_fn + = builtin_decl_explicit (BUILT_IN_OMP_INIT_ALLOCATOR); + tree omp_destroy_allocator_fn + = builtin_decl_explicit (BUILT_IN_OMP_DESTROY_ALLOCATOR); + + for (tree *cp = &OMP_CLAUSES (expr); *cp != NULL;) + if (OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_USES_ALLOCATORS) + { + tree c = *cp; + tree allocator = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c); + tree memspace = OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c); + tree traits = OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c); + tree ntraits + = ((traits + && DECL_INITIAL (traits) + && TREE_CODE (DECL_INITIAL (traits)) == CONSTRUCTOR) + ? build_int_cst (integer_type_node, + CONSTRUCTOR_NELTS (DECL_INITIAL (traits))) + : integer_zero_node); + tree traits_var + = (traits != NULL_TREE + ? get_initialized_tmp_var (DECL_INITIAL (traits), + &init_seq, NULL) + : null_pointer_node); + + tree memspace_var = create_tmp_var (pointer_sized_int_node, + "memspace_enum"); + if (memspace == NULL_TREE) + memspace = build_int_cst (pointer_sized_int_node, 0); + else + memspace = fold_convert (pointer_sized_int_node, + memspace); + g = gimple_build_assign (memspace_var, memspace); + gimple_seq_add_stmt (&init_seq, g); + + tree initcall = build_call_expr_loc (OMP_CLAUSE_LOCATION (c), + omp_init_allocator_fn, 3, + memspace_var, + ntraits, + traits_var); + initcall = fold_convert (TREE_TYPE (allocator), initcall); + gimplify_assign (allocator, initcall, &init_seq); + + g = gimple_build_call (omp_destroy_allocator_fn, 1, allocator); + gimple_seq_add_stmt (&fini_seq, g); + + /* Finished generating runtime calls, remove USES_ALLOCATORS + clause. */ + *cp = OMP_CLAUSE_CHAIN (c); + } + else + cp = &OMP_CLAUSE_CHAIN (*cp); + + if (fini_seq) + { + gbind *bind = as_a (gimple_seq_first_stmt (body)); + g = gimple_build_try (gimple_bind_body (bind), + fini_seq, GIMPLE_TRY_FINALLY); + gimple_seq_add_stmt (&init_seq, g); + gimple_bind_set_body (bind, init_seq); + } + } } else gimplify_and_add (OMP_BODY (expr), &body); diff --git a/gcc/omp-builtins.def b/gcc/omp-builtins.def index e442b0b5c94..63482c1232d 100644 --- a/gcc/omp-builtins.def +++ b/gcc/omp-builtins.def @@ -66,6 +66,10 @@ DEF_GOMP_BUILTIN (BUILT_IN_OMP_GET_TEAM_NUM, "omp_get_team_num", BT_FN_INT, ATTR_CONST_NOTHROW_LEAF_LIST) DEF_GOMP_BUILTIN (BUILT_IN_OMP_GET_NUM_TEAMS, "omp_get_num_teams", BT_FN_INT, ATTR_CONST_NOTHROW_LEAF_LIST) +DEF_GOMP_BUILTIN (BUILT_IN_OMP_INIT_ALLOCATOR, "omp_init_allocator", + BT_FN_PTRMODE_PTRMODE_INT_PTR, ATTR_NOTHROW_LEAF_LIST) +DEF_GOMP_BUILTIN (BUILT_IN_OMP_DESTROY_ALLOCATOR, "omp_destroy_allocator", + BT_FN_VOID_PTRMODE, ATTR_NOTHROW_LEAF_LIST) DEF_GOMP_BUILTIN (BUILT_IN_GOMP_ATOMIC_START, "GOMP_atomic_start", BT_FN_VOID, ATTR_NOTHROW_LEAF_LIST) diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 68794c01c43..77779f48cc0 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -5291,7 +5291,12 @@ lower_private_allocate (tree var, tree new_var, tree &allocator, allocator = TREE_PURPOSE (allocator); } if (TREE_CODE (allocator) != INTEGER_CST) - allocator = build_outer_var_ref (allocator, ctx); + { + if (is_task_ctx (ctx)) + allocator = build_receiver_ref (allocator, false, ctx); + else + allocator = build_outer_var_ref (allocator, ctx); + } allocator = fold_convert (pointer_sized_int_node, allocator); if (TREE_CODE (allocator) != INTEGER_CST) { @@ -6275,7 +6280,7 @@ lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist, if (TREE_CODE (allocator) == TREE_LIST) allocator = TREE_PURPOSE (allocator); if (TREE_CODE (allocator) != INTEGER_CST) - allocator = build_outer_var_ref (allocator, ctx); + allocator = build_receiver_ref (allocator, false, ctx); allocator = fold_convert (pointer_sized_int_node, allocator); allocate_ptr = unshare_expr (x); @@ -6595,7 +6600,7 @@ lower_rec_input_clauses (tree clauses, gimple_seq *ilist, gimple_seq *dlist, if (TREE_CODE (allocator) == TREE_LIST) allocator = TREE_PURPOSE (allocator); if (TREE_CODE (allocator) != INTEGER_CST) - allocator = build_outer_var_ref (allocator, ctx); + allocator = build_receiver_ref (allocator, false, ctx); allocator = fold_convert (pointer_sized_int_node, allocator); allocate_ptr = unshare_expr (x); @@ -12789,6 +12794,8 @@ create_task_copyfn (gomp_task *task_stmt, omp_context *ctx) allocator = *tcctx.cb.decl_map->get (allocator); tree a = build_simple_mem_ref_loc (loc, sarg); allocator = omp_build_component_ref (a, allocator); + if (POINTER_TYPE_P (TREE_TYPE (allocator))) + allocator = build_simple_mem_ref (allocator); } allocator = fold_convert (pointer_sized_int_node, allocator); tree a = builtin_decl_explicit (BUILT_IN_GOMP_ALLOC); diff --git a/gcc/system.h b/gcc/system.h index ce3bcbba753..9899b956ef0 100644 --- a/gcc/system.h +++ b/gcc/system.h @@ -1292,6 +1292,14 @@ void gcc_stablesort_r (void *, size_t, size_t, sort_r_cmp_fn *, void *data); #define NULL nullptr #endif +/* Return true if STR string starts with PREFIX. */ + +static inline bool +startswith (const char *str, const char *prefix) +{ + return strncmp (str, prefix, strlen (prefix)) == 0; +} + /* Return true if STR string ends with SUFFIX. */ static inline bool diff --git a/gcc/testsuite/c-c++-common/gomp/uses_allocators-1.c b/gcc/testsuite/c-c++-common/gomp/uses_allocators-1.c new file mode 100644 index 00000000000..29541abd525 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/uses_allocators-1.c @@ -0,0 +1,46 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" } */ + +#include + +omp_alloctrait_key_t k; +omp_alloctrait_value_t v; + +int main (void) +{ + omp_allocator_handle_t foo, bar; + const omp_alloctrait_t foo_traits[] = { { omp_atk_pinned, omp_atv_true }, + { omp_atk_partition, omp_atv_nearest } }; + #pragma omp target + ; + #pragma omp target uses_allocators (bar) + ; + #pragma omp target uses_allocators (foo (foo_traits)) + ; + #pragma omp target uses_allocators (foo (foo_traits), bar (foo_traits)) + ; + #pragma omp target uses_allocators (memspace(omp_high_bw_mem_space) : foo) + ; + #pragma omp target uses_allocators (traits(foo_traits) : bar) + ; + #pragma omp target parallel uses_allocators (memspace(omp_high_bw_mem_space), traits(foo_traits) : bar) + ; + #pragma omp target parallel uses_allocators (traits(foo_traits), memspace(omp_high_bw_mem_space) : bar) uses_allocators(foo) + { + void *p = omp_alloc ((unsigned long) 32, bar); + omp_free (p, bar); + } + return 0; +} + +/* { dg-final { scan-tree-dump "pragma omp target" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target private\\(bar\\) uses_allocators\\(bar: memspace\\(\\), traits\\(\\)\\)" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target private\\(foo\\) uses_allocators\\(foo: memspace\\(\\), traits\\(foo_traits\\)\\)" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target private\\(bar\\) uses_allocators\\(bar: memspace\\(\\), traits\\(foo_traits\\)\\) private\\(foo\\) uses_allocators\\(foo: memspace\\(\\), traits\\(foo_traits\\)\\)" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target private\\(foo\\) uses_allocators\\(foo: memspace\\(omp_high_bw_mem_space\\), traits\\(\\)\\)" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target private\\(bar\\) uses_allocators\\(bar: memspace\\(\\), traits\\(foo_traits\\)\\)" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target private\\(bar\\) uses_allocators\\(bar: memspace\\(omp_high_bw_mem_space\\), traits\\(foo_traits\\)\\)" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target private\\(bar\\) uses_allocators\\(bar: memspace\\(omp_high_bw_mem_space\\), traits\\(foo_traits\\)\\) private\\(foo\\) uses_allocators\\(foo: memspace\\(\\), traits\\(\\)\\)" "original" } } */ + +/* { dg-final { scan-tree-dump-times "__builtin_omp_init_allocator" 9 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "__builtin_omp_destroy_allocator" 9 "gimple" } } */ diff --git a/gcc/testsuite/c-c++-common/gomp/uses_allocators-2.c b/gcc/testsuite/c-c++-common/gomp/uses_allocators-2.c new file mode 100644 index 00000000000..f350c0a409e --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/uses_allocators-2.c @@ -0,0 +1,39 @@ +/* { dg-do compile } */ + +#include + +omp_alloctrait_key_t k; +omp_alloctrait_value_t v; + +int main (void) +{ + omp_allocator_handle_t foo, bar; + const omp_alloctrait_t traits_array[] = { { omp_atk_pinned, omp_atv_true }, + { omp_atk_partition, omp_atv_nearest } }; + + #pragma omp target uses_allocators (baz) /* { dg-error "'baz' undeclared .first use in this function." "" { target c } } */ + ; /* { dg-error "'baz' has not been declared" "" { target c++ } .-1 } */ + #pragma omp target uses_allocators (foo (xyz)) /* { dg-error "'xyz' undeclared .first use in this function." "" { target c } } */ + ; /* { dg-error "'xyz' has not been declared" "" { target c++ } .-1 } */ + #pragma omp target uses_allocators (foo (traits_array), baz (traits_array)) /* { dg-error "'baz' has not been declared" "" { target c++ } } */ + ; + #pragma omp target uses_allocators (memspace(omp_no_such_space) : foo) /* { dg-error "'omp_no_such_space' undeclared .first use in this function." "" { target c } } */ + ; /* { dg-error "'omp_no_such_space' has not been declared" "" { target c++ } .-1 } */ + #pragma omp target uses_allocators (memspace(1) : foo) /* { dg-error "expected identifier before numeric constant" } */ + ; /* { dg-error "expected '\\\)' before ':' token" "" { target c } .-1 } */ + #pragma omp target uses_allocators (memspace(omp_no_such_space) : foo, bar) /* { dg-error "'uses_allocators' clause only accepts a single allocator when using modifiers" } */ + ; /* { dg-error "'omp_no_such_space' has not been declared" "" { target c++ } .-1 } */ + #pragma omp target uses_allocators (traits(xyz) : bar) /* { dg-error "traits array must be of 'const omp_alloctrait_t \\\[\\\]' type" "" { target c } } */ + ; /* { dg-error "'xyz' has not been declared" "" { target c++ } .-1 } */ + #pragma omp target uses_allocators (memspace(omp_high_bw_mem_space), traits(traits_array), memspace (omp_no_such_space) : bar) /* { dg-error "duplicate 'memspace' modifier" } */ + ; + #pragma omp target uses_allocators (traitz(traits_array), memspace(omp_high_bw_mem_space) : bar) /* { dg-error "unknown modifier 'traitz'" } */ + ; + #pragma omp target uses_allocators (omp_null_allocator) /* { dg-error "'omp_null_allocator' cannot be used in 'uses_allocators' clause" } */ + ; + #pragma omp target uses_allocators (memspace(omp_high_bw_mem_space) : foo, bar) /* { dg-error "'uses_allocators' clause only accepts a single allocator when using modifiers" } */ + ; + #pragma omp target uses_allocators (memspace(omp_high_bw_mem_space) : foo(foo_traits)) /* { dg-error "legacy 'foo\\\(foo_traits\\\)' traits syntax not allowed in 'uses_allocators' clause when using modifiers" } */ + ; + return 0; +} diff --git a/gcc/testsuite/c-c++-common/gomp/uses_allocators-3.c b/gcc/testsuite/c-c++-common/gomp/uses_allocators-3.c new file mode 100644 index 00000000000..80b2844729a --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/uses_allocators-3.c @@ -0,0 +1,31 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" } */ + +#include + +int main (void) +{ + omp_allocator_handle_t memspace, traits; + const omp_alloctrait_t mytraits[] = { { omp_atk_pinned, omp_atv_true }, + { omp_atk_partition, omp_atv_nearest } }; + #pragma omp target uses_allocators (memspace) + ; + #pragma omp target uses_allocators (traits) + ; + #pragma omp target uses_allocators (traits, memspace) + ; + #pragma omp target uses_allocators (traits (mytraits)) + ; + #pragma omp target uses_allocators (memspace (mytraits), omp_default_mem_alloc) + ; + return 0; +} + +/* { dg-final { scan-tree-dump "pragma omp target private\\(memspace\\) uses_allocators\\(memspace: memspace\\(\\), traits\\(\\)\\)" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target private\\(traits\\) uses_allocators\\(traits: memspace\\(\\), traits\\(\\)\\)" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target private\\(memspace\\) uses_allocators\\(memspace: memspace\\(\\), traits\\(\\)\\) private\\(traits\\) uses_allocators\\(traits: memspace\\(\\), traits\\(\\)\\)" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target private\\(traits\\) uses_allocators\\(traits: memspace\\(\\), traits\\(mytraits\\)\\)" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target private\\(memspace\\) uses_allocators\\(memspace: memspace\\(\\), traits\\(mytraits\\)\\)" "original" } } */ + +/* { dg-final { scan-tree-dump-times "__builtin_omp_init_allocator" 6 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "__builtin_omp_destroy_allocator" 6 "gimple" } } */ diff --git a/gcc/testsuite/gfortran.dg/gomp/allocate-1.f90 b/gcc/testsuite/gfortran.dg/gomp/allocate-1.f90 index 8bc6b768778..f5707899eff 100644 --- a/gcc/testsuite/gfortran.dg/gomp/allocate-1.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/allocate-1.f90 @@ -80,7 +80,8 @@ subroutine foo(x, y) !$omp target teams distribute parallel do private (x) firstprivate (y) & !$omp allocate ((omp_default_mem_alloc + 0):z) allocate & - !$omp (omp_default_mem_alloc: x, y) allocate (h: r) lastprivate (z) reduction(+:r) + !$omp (omp_default_mem_alloc: x, y) allocate (h: r) lastprivate (z) reduction(+:r) & + !$omp uses_allocators (h) do i = 1, 10 call bar (0, x, z); call bar2 (1, y, r); diff --git a/gcc/testsuite/gfortran.dg/gomp/uses_allocators-1.f90 b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-1.f90 new file mode 100644 index 00000000000..4ca76e7004c --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-1.f90 @@ -0,0 +1,53 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" } + +program main + use omp_lib + implicit none + integer, allocatable :: arr(:) + integer (omp_allocator_handle_kind) :: bar, foo + + type (omp_alloctrait), parameter :: traits_array(*) = & + [omp_alloctrait(omp_atk_pinned,omp_atv_true),& + omp_alloctrait(omp_atk_partition,omp_atv_nearest)] + + !$omp target allocate(bar : arr) uses_allocators(bar) + block + allocate(arr(100)) + end block + + !$omp target uses_allocators(omp_default_mem_alloc) + block + end block + + !$omp target uses_allocators(bar(traits_array), foo (traits_array)) + block + end block + + !$omp target uses_allocators(traits(traits_array) : bar) + block + end block + + !$omp target parallel uses_allocators(memspace (omp_low_lat_mem_space) : bar) + block + end block + + !$omp target parallel uses_allocators(memspace (omp_high_bw_mem_space), traits(traits_array) : bar) + block + use iso_c_binding + type(c_ptr) :: ptr + integer(c_size_t) :: sz = 32 + ptr = omp_alloc (sz, bar) + call omp_free (ptr, bar) + end block + +end program main + +! { dg-final { scan-tree-dump "pragma omp target allocate\\(allocator\\(bar\\):arr\\) private\\(bar\\) uses_allocators\\(bar: memspace\\(\\), traits\\(\\)\\)" "original" } } +! { dg-final { scan-tree-dump "pragma omp target" "original" } } +! { dg-final { scan-tree-dump "pragma omp target private\\(foo\\) uses_allocators\\(foo: memspace\\(\\), traits\\(traits_array\\)\\) private\\(bar\\) uses_allocators\\(bar: memspace\\(\\), traits\\(traits_array\\)\\)" "original" } } +! { dg-final { scan-tree-dump "pragma omp target private\\(bar\\) uses_allocators\\(bar: memspace\\(\\), traits\\(traits_array\\)\\)" "original" } } +! { dg-final { scan-tree-dump "pragma omp target private\\(bar\\) uses_allocators\\(bar: memspace\\(.\\), traits\\(\\)\\)" "original" } } +! { dg-final { scan-tree-dump "pragma omp target private\\(bar\\) uses_allocators\\(bar: memspace\\(.\\), traits\\(traits_array\\)\\)" "original" } } +! { dg-final { scan-tree-dump-times "__builtin_omp_init_allocator" 6 "gimple" } } +! { dg-final { scan-tree-dump-times "__builtin_omp_destroy_allocator" 6 "gimple" } } diff --git a/gcc/testsuite/gfortran.dg/gomp/uses_allocators-2.f90 b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-2.f90 new file mode 100644 index 00000000000..ce5e8b3298b --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-2.f90 @@ -0,0 +1,52 @@ +! { dg-do compile } + +program main + use omp_lib + implicit none + integer (omp_allocator_handle_kind) :: bar, foo + + type (omp_alloctrait), parameter :: traits_array(*) = & + [omp_alloctrait(omp_atk_pinned,omp_atv_true),& + omp_alloctrait(omp_atk_partition,omp_atv_nearest)] + + !$omp target uses_allocators(omp_non_existant_alloc) ! { dg-error "Symbol 'omp_non_existant_alloc' at .1. is ambiguous" } + block + end block + + !$omp target uses_allocators(bar(traits_array), foo (traits_array), ) ! { dg-error "Expected identifier at .1." } + block ! { dg-error "Failed to match clause at .1." "" { target *-*-* } .-1 } + end block + + !$omp target uses_allocators(traits(xyz) : bar) ! { dg-error "Symbol 'xyz' at .1. is ambiguous" } + block + end block + + !$omp target uses_allocators(memspace(omp_non_existant_mem_space) : foo) ! { dg-error "Symbol 'omp_non_existant_mem_space' at .1. is ambiguous" } + block + end block + + !$omp target uses_allocators(traits(traits_array), traits(traits_array) : bar) ! { dg-error "duplicate 'traits' modifier at .1." } + block + end block + + !$omp target uses_allocators(memspace(omp_default_mem_space), memspace(omp_default_mem_space) : foo) ! { dg-error "duplicate 'memspace' modifier at .1." } + block + end block + + !$omp target uses_allocators(memspace(omp_default_mem_space), traits(traits_array), traits(traits_array) : foo) ! { dg-error "duplicate 'traits' modifier at .1." } + block + end block + + !$omp target uses_allocators (omp_null_allocator) ! { dg-error "'omp_null_allocator' cannot be used in 'uses_allocators' clause at .1." } + block + end block + + !$omp target uses_allocators (memspace(omp_high_bw_mem_space) : foo, bar) ! { dg-error "'uses_allocators' clause only accepts a single allocator when using modifiers at .1." } + block + end block + + !$omp target uses_allocators (memspace(omp_high_bw_mem_space) : foo(foo_traits)) ! { dg-error "legacy 'foo\\\(foo_traits\\\)' traits syntax not allowed in 'uses_allocators' clause when using modifiers at .1." } + block + end block + +end program main diff --git a/gcc/testsuite/gfortran.dg/gomp/uses_allocators-3.f90 b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-3.f90 new file mode 100644 index 00000000000..0f024264700 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-3.f90 @@ -0,0 +1,14 @@ +! { dg-do compile } + +program main + use omp_lib + implicit none + integer, allocatable :: arr(:) + integer (omp_allocator_handle_kind) :: bar + + !$omp target allocate(bar : arr) ! { dg-error "allocator ''bar'' requires 'uses_allocators.bar.' clause in target region" } + block + allocate(arr(100)) + end block + +end program main diff --git a/gcc/tree-core.h b/gcc/tree-core.h index 7cb9c51ccd4..0e56706f2b7 100644 --- a/gcc/tree-core.h +++ b/gcc/tree-core.h @@ -516,7 +516,10 @@ enum omp_clause_code { OMP_CLAUSE_FINALIZE, /* OpenMP clause: allocator. */ - OMP_CLAUSE_ALLOCATOR + OMP_CLAUSE_ALLOCATOR, + + /* OpenMP clause: uses_allocators. */ + OMP_CLAUSE_USES_ALLOCATORS, }; #undef DEFTREESTRUCT diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c index fd336d3a216..36ccf899476 100644 --- a/gcc/tree-pretty-print.c +++ b/gcc/tree-pretty-print.c @@ -771,6 +771,20 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) pp_right_paren (pp); break; + case OMP_CLAUSE_USES_ALLOCATORS: + pp_string (pp, "uses_allocators("); + dump_generic_node (pp, OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (clause), + spc, flags, false); + pp_string (pp, ": memspace("); + dump_generic_node (pp, OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (clause), + spc, flags, false); + pp_string (pp, "), traits("); + dump_generic_node (pp, OMP_CLAUSE_USES_ALLOCATORS_TRAITS (clause), + spc, flags, false); + pp_right_paren (pp); + pp_right_paren (pp); + break; + case OMP_CLAUSE_AFFINITY: pp_string (pp, "affinity("); { diff --git a/gcc/tree.c b/gcc/tree.c index 1b99620bfcb..62a17f2c549 100644 --- a/gcc/tree.c +++ b/gcc/tree.c @@ -365,6 +365,7 @@ unsigned const char omp_clause_num_ops[] = 0, /* OMP_CLAUSE_IF_PRESENT */ 0, /* OMP_CLAUSE_FINALIZE */ 2, /* OMP_CLAUSE_ALLOCATOR */ + 3, /* OMP_CLAUSE_USES_ALLOCATORS */ }; const char * const omp_clause_code_name[] = @@ -454,6 +455,8 @@ const char * const omp_clause_code_name[] = "tile", "if_present", "finalize", + "allocator", + "uses_allocators", }; diff --git a/gcc/tree.h b/gcc/tree.h index 414520696cf..5caea7c9dea 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -1833,6 +1833,15 @@ class auto_suppress_location_wrappers #define OMP_CLAUSE_ALLOCATE_COMBINED(NODE) \ (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_ALLOCATE)->base.public_flag) +#define OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR(NODE) \ + OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_USES_ALLOCATORS), 0) + +#define OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE(NODE) \ + OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_USES_ALLOCATORS), 1) + +#define OMP_CLAUSE_USES_ALLOCATORS_TRAITS(NODE) \ + OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_USES_ALLOCATORS), 2) + #define OMP_CLAUSE_NUM_TEAMS_UPPER_EXPR(NODE) \ OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_NUM_TEAMS), 0)