From patchwork Fri May 6 13:20:48 2022 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Chung-Lin Tang X-Patchwork-Id: 53551 Return-Path: X-Original-To: patchwork@sourceware.org Delivered-To: patchwork@sourceware.org Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 74423395445B for ; Fri, 6 May 2022 13:21:54 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa1.mentor.iphmx.com (esa1.mentor.iphmx.com [68.232.129.153]) by sourceware.org (Postfix) with ESMTPS id 462CA394FC0E for ; Fri, 6 May 2022 13:20:59 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 462CA394FC0E Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com X-IronPort-AV: E=Sophos;i="5.91,203,1647331200"; d="scan'208";a="78055070" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa1.mentor.iphmx.com with ESMTP; 06 May 2022 05:20:57 -0800 IronPort-SDR: P6ho2BbVBLoKShPxej0AW867VRMyYEXykmcGF4L3bGvR8+2XR5Wq0yJsgbX5HaJXOdQiKFcXgs HC/nZvxKlapNDyE/Nl7x0N2BJYf+zsm4oDJasFoBSx/6kpud/Ofw/qtwgOsV2GIpq0TimybHXc rXywCeyDhUh22DFleoF31Bnxlg+Qq37l1LDonBoQd+DkfoLiF4U+oKAp/f7qye2b8D5cgNE+VS 6i70BLvvomxLkd+e/HHztsvxeL0zdRsRzu2rK5pyBUVJjef5anP43W0OP3ysaqZT5gl2u9m99b XMw= Message-ID: <46d77e14-080c-db6c-4032-e12899c5d059@codesourcery.com> Date: Fri, 6 May 2022 21:20:48 +0800 MIME-Version: 1.0 User-Agent: Mozilla/5.0 (Macintosh; Intel Mac OS X 10.13; rv:91.0) Gecko/20100101 Thunderbird/91.9.0 Content-Language: en-US From: Chung-Lin Tang To: gcc-patches , Jakub Jelinek , Tobias Burnus , Catherine Moore , Andrew Stubbs , Hafiz Abid Qadeer Subject: [PATCH, OpenMP] Implement uses_allocators clause for target regions X-ClientProxiedBy: svr-orw-mbx-08.mgc.mentorg.com (147.34.90.208) To svr-orw-mbx-02.mgc.mentorg.com (147.34.90.202) X-Spam-Status: No, score=-10.4 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, KAM_MANYTO, KAM_SHORT, SPF_HELO_PASS, SPF_PASS, TXREP, T_SCC_BODY_TEXT_LINE autolearn=ham autolearn_force=no version=3.4.4 X-Spam-Checker-Version: SpamAssassin 3.4.4 (2020-01-24) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces+patchwork=sourceware.org@gcc.gnu.org Sender: "Gcc-patches" Hi Jakub, this patch implements the uses_allocators clause for OpenMP target regions. 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", as to mentioned in[1]. This is done during gimplify_scan_omp_clauses. [1] https://gcc.gnu.org/pipermail/gcc-patches/2022-May/594039.html Tested on mainline, please see if this is okay. Thanks, Chung-Lin 2022-05-06 Chung-Lin Tang 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[]. * 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. gcc/testsuite/ChangeLog: * c-c++-common/gomp/uses_allocators-1.c: New test. * c-c++-common/gomp/uses_allocators-2.c: New test. * 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 --git a/gcc/builtin-types.def b/gcc/builtin-types.def index 3a7cecdf087..be3e6ff697e 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) @@ -641,6 +642,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.cc b/gcc/c-family/c-omp.cc index 777cdc65572..5066e137cf4 100644 --- a/gcc/c-family/c-omp.cc +++ b/gcc/c-family/c-omp.cc @@ -1870,6 +1870,7 @@ c_omp_split_clauses (location_t loc, enum tree_code code, case OMP_CLAUSE_HAS_DEVICE_ADDR: 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 54864c2ec41..7f8944f81d6 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.cc b/gcc/c/c-parser.cc index 129dd727ef3..2befad91158 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -12907,6 +12907,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)) @@ -15624,6 +15626,225 @@ c_parser_omp_clause_allocate (c_parser *parser, tree list) return nl; } +/* OpenMP 5.2: + uses_allocators ( allocator-list ) + + allocator-list: + allocator + allocator , allocator-list + allocator ( traits-array ) + allocator ( traits-array ) , allocator-list + + 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 nl; + matching_parens parens; + if (!parens.require_open (parser)) + return list; + + bool has_modifiers = false; + tree memspace_expr = NULL_TREE; + tree traits_var = NULL_TREE; + + if (c_parser_next_token_is (parser, CPP_NAME)) + { + c_token *tok = c_parser_peek_token (parser); + const char *p = IDENTIFIER_POINTER (tok->value); + + if (strcmp ("traits", p) == 0 || strcmp ("memspace", p) == 0) + { + has_modifiers = true; + c_parser_consume_token (parser); + matching_parens parens2;; + parens2.require_open (parser); + + if (c_parser_next_token_is (parser, CPP_NAME) + && (c_parser_peek_token (parser)->id_kind == C_ID_ID + || c_parser_peek_token (parser)->id_kind == C_ID_TYPENAME)) + { + tok = c_parser_peek_token (parser); + tree t = lookup_name (tok->value); + + if (t == NULL_TREE) + { + undeclared_variable (tok->location, tok->value); + t = error_mark_node; + } + else + { + if (strcmp ("memspace", p) == 0) + memspace_expr = t; + else + traits_var = t; + } + c_parser_consume_token (parser); + } + + parens2.skip_until_found_close (parser); + + if (c_parser_next_token_is (parser, CPP_COMMA)) + { + c_parser_consume_token (parser); + tok = c_parser_peek_token (parser); + const char *q = ""; + if (c_parser_next_token_is (parser, CPP_NAME)) + q = IDENTIFIER_POINTER (tok->value); + if (strcmp (q, "memspace") != 0 && strcmp (q, "traits") != 0) + { + c_parser_error (parser, "expected % or %"); + parens.skip_until_found_close (parser); + return list; + } + else if (strcmp (p, q) == 0) + { + error_at (tok->location, "duplicate %qs modifier", p); + parens.skip_until_found_close (parser); + return list; + } + c_parser_consume_token (parser); + if (!parens2.require_open (parser)) + { + parens.skip_until_found_close (parser); + return list; + } + + if (c_parser_next_token_is (parser, CPP_NAME) + && (c_parser_peek_token (parser)->id_kind == C_ID_ID + || c_parser_peek_token (parser)->id_kind == C_ID_TYPENAME)) + { + tok = c_parser_peek_token (parser); + tree t = lookup_name (tok->value); + if (t == NULL_TREE) + { + undeclared_variable (tok->location, tok->value); + t = error_mark_node; + } + else + { + if (strcmp ("memspace", q) == 0) + memspace_expr = t; + else + traits_var = t; + } + c_parser_consume_token (parser); + } + parens2.skip_until_found_close (parser); + } + } + } + + if (has_modifiers) + { + if (!c_parser_require (parser, CPP_COLON, "expected %<:%>")) + { + parens.skip_until_found_close (parser); + return list; + } + + if (c_parser_next_token_is (parser, CPP_NAME) + && c_parser_peek_token (parser)->id_kind == C_ID_ID) + { + tree t = lookup_name (c_parser_peek_token (parser)->value); + + if (t == NULL_TREE) + { + undeclared_variable (c_parser_peek_token (parser)->location, + c_parser_peek_token (parser)->value); + t = error_mark_node; + } + else if (t != error_mark_node) + { + 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) = traits_var; + OMP_CLAUSE_CHAIN (c) = list; + + nl = c; + } + c_parser_consume_token (parser); + + if (c_parser_next_token_is (parser, CPP_OPEN_PAREN)) + c_parser_error (parser, "modifiers cannot be used with " + "legacy array syntax"); + } + else + c_parser_error (parser, "expected identifier"); + } + else + { + traits_var = NULL_TREE; + + nl = list; + while (true) + { + if (c_parser_next_token_is (parser, CPP_NAME) + && c_parser_peek_token (parser)->id_kind == C_ID_ID) + { + c_token *tok = c_parser_peek_token (parser); + tree t = lookup_name (tok->value); + + if (t == NULL_TREE) + { + undeclared_variable (tok->location, tok->value); + break; + } + else if (t == error_mark_node) + break; + + 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) + && c_parser_peek_token (parser)->id_kind == C_ID_ID) + { + tok = c_parser_peek_token (parser); + traits_var = lookup_name (tok->value); + if (traits_var == NULL_TREE) + { + undeclared_variable (tok->location, tok->value); + break; + } + c_parser_consume_token (parser); + } + else + c_parser_error (parser, "expected identifier"); + parens2.require_close (parser); + } + + tree c = build_omp_clause (clause_loc, + OMP_CLAUSE_USES_ALLOCATORS); + OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c) = t; + OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c) = NULL_TREE; + OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c) = traits_var; + OMP_CLAUSE_CHAIN (c) = nl; + + nl = c; + } + + if (c_parser_next_token_is_not (parser, CPP_COMMA)) + break; + c_parser_consume_token (parser); + } + } + + parens.skip_until_found_close (parser); + return nl; +} + /* OpenMP 4.0: linear ( variable-list ) linear ( variable-list : expression ) @@ -17050,6 +17271,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"; @@ -21061,7 +21286,8 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser, | (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_HAS_DEVICE_ADDR)) + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR)\ + | (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.cc b/gcc/c/c-typeck.cc index e130196a3a7..0e1f33b655d 100644 --- a/gcc/c/c-typeck.cc +++ b/gcc/c/c-typeck.cc @@ -14763,6 +14763,102 @@ 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 (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; + } + 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; + } + if (TREE_CODE (t) == CONST_DECL) + { + 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"); + + /* Currently for pre-defined allocators in libgomp, we do not + require additional init/fini inside target regions, so discard + such clauses. */ + remove = true; + } + 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; + } + t = OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c); + if (t != NULL_TREE) + { + bool type_err = false; + + if (TREE_CODE (TREE_TYPE (t)) != ARRAY_TYPE) + 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.cc b/gcc/cp/parser.cc index 2235da10c7c..143c70d0b49 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -36490,6 +36490,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)) @@ -38733,6 +38735,246 @@ cp_parser_omp_clause_allocate (cp_parser *parser, tree list) return nlist; } +/* OpenMP 5.2: + uses_allocators ( allocator-list ) + + allocator-list: + allocator + allocator , allocator-list + allocator ( traits-array ) + allocator ( traits-array ) , allocator-list + + 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; + matching_parens parens; + if (!parens.require_open (parser)) + return list; + + bool has_modifiers = false; + tree memspace_expr = NULL_TREE; + tree traits_var = NULL_TREE; + + if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)) + { + cp_token *tok = cp_lexer_peek_token (parser->lexer); + const char *p = IDENTIFIER_POINTER (tok->u.value); + + if (strcmp ("traits", p) == 0 || strcmp ("memspace", p) == 0) + { + has_modifiers = true; + cp_lexer_consume_token (parser->lexer); + matching_parens parens2;; + parens2.require_open (parser); + + if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)) + { + tok = cp_lexer_peek_token (parser->lexer); + tree id = tok->u.value; + + t = cp_parser_lookup_name_simple (parser, id, tok->location); + if (t == error_mark_node) + cp_parser_name_lookup_error (parser, id, t, NLE_NULL, + tok->location); + else + { + if (strcmp ("memspace", p) == 0) + memspace_expr = t; + else + traits_var = t; + } + cp_lexer_consume_token (parser->lexer); + } + + if (t == error_mark_node || !parens.require_close (parser)) + { + cp_parser_skip_to_closing_parenthesis (parser, + /*recovering=*/true, + /*or_comma=*/false, + /*consume_paren=*/true); + return list; + } + + if (cp_lexer_next_token_is (parser->lexer, CPP_COMMA)) + { + cp_lexer_consume_token (parser->lexer); + tok = cp_lexer_peek_token (parser->lexer); + const char *q = ""; + + if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)) + q = IDENTIFIER_POINTER (tok->u.value); + + if (strcmp (q, "memspace") != 0 && strcmp (q, "traits") != 0) + { + cp_parser_error (parser, "expected % or %"); + cp_parser_skip_to_closing_parenthesis (parser, + /*recovering=*/true, + /*or_comma=*/false, + /*consume_paren=*/true); + return list; + } + else if (strcmp (p, q) == 0) + { + error_at (tok->location, "duplicate %qs modifier", p); + cp_parser_skip_to_closing_parenthesis (parser, + /*recovering=*/true, + /*or_comma=*/false, + /*consume_paren=*/true); + return list; + } + cp_lexer_consume_token (parser->lexer); + if (!parens2.require_open (parser)) + { + cp_parser_skip_to_closing_parenthesis (parser, + /*recovering=*/true, + /*or_comma=*/false, + /*consume_paren=*/true); + return list; + } + + if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)) + { + tok = cp_lexer_peek_token (parser->lexer); + tree id = tok->u.value; + + t = cp_parser_lookup_name_simple (parser, id, tok->location); + if (t == error_mark_node) + cp_parser_name_lookup_error (parser, id, t, NLE_NULL, + tok->location); + else + { + if (strcmp ("memspace", q) == 0) + memspace_expr = t; + else + traits_var = t; + } + cp_lexer_consume_token (parser->lexer); + } + + if (t == error_mark_node || !parens.require_close (parser)) + { + cp_parser_skip_to_closing_parenthesis (parser, + /*recovering=*/true, + /*or_comma=*/false, + /*consume_paren=*/true); + return list; + } + } + + } + } + + if (has_modifiers) + { + if (!cp_parser_require (parser, CPP_COLON, RT_COLON)) + { + cp_parser_skip_to_closing_parenthesis (parser, /*recovering=*/true, + /*or_comma=*/false, + /*consume_paren=*/true); + return list; + } + + if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)) + { + cp_token *tok = cp_lexer_peek_token (parser->lexer); + tree id = tok->u.value; + tree t = cp_parser_lookup_name_simple (parser, id, tok->location); + + if (t == error_mark_node) + cp_parser_name_lookup_error (parser, id, t, NLE_NULL, + tok->location); + else + { + 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) = traits_var; + OMP_CLAUSE_CHAIN (c) = list; + + nl = c; + } + cp_lexer_consume_token (parser->lexer); + + if (cp_lexer_next_token_is (parser->lexer, CPP_OPEN_PAREN)) + cp_parser_error (parser, "modifiers cannot be used with " + "legacy array syntax"); + } + else + cp_parser_error (parser, "expected identifier"); + } + else + { + traits_var = NULL_TREE; + + nl = list; + while (true) + { + if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)) + { + cp_token *tok = cp_lexer_peek_token (parser->lexer); + tree id = tok->u.value; + + tree t = cp_parser_lookup_name_simple (parser, id, tok->location); + + if (t == error_mark_node) + cp_parser_name_lookup_error (parser, id, t, NLE_NULL, + 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); + id = tok->u.value; + traits_var = cp_parser_lookup_name_simple (parser, id, + tok->location); + if (traits_var == error_mark_node) + { + cp_parser_name_lookup_error (parser, id, traits_var, + NLE_NULL, tok->location); + break; + } + cp_lexer_consume_token (parser->lexer); + } + else + cp_parser_error (parser, "expected identifier"); + parens2.require_close (parser); + } + + tree c = build_omp_clause (clause_loc, + OMP_CLAUSE_USES_ALLOCATORS); + OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c) = t; + OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c) = NULL_TREE; + OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c) = traits_var; + OMP_CLAUSE_CHAIN (c) = nl; + + nl = c; + } + + if (cp_lexer_next_token_is_not (parser->lexer, CPP_COMMA)) + break; + cp_lexer_consume_token (parser->lexer); + } + } + + parens.require_close (parser); + return nl; +} + /* OpenMP 2.5: lastprivate ( variable-list ) @@ -40283,6 +40525,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; @@ -44291,7 +44537,8 @@ cp_parser_omp_target_update (cp_parser *parser, cp_token *pragma_tok, | (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_HAS_DEVICE_ADDR)) + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR)\ + | (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.cc b/gcc/cp/semantics.cc index 377f61113c0..ac8c4b2e3e9 100644 --- a/gcc/cp/semantics.cc +++ b/gcc/cp/semantics.cc @@ -7772,6 +7772,90 @@ 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 (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; + } + if (TREE_CODE (t) == CONST_DECL) + { + 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"); + + /* Currently for pre-defined allocators in libgomp, we do not + require additional init/fini inside target regions, so discard + such clauses. */ + remove = true; + } + 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; + } + t = OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c); + if (t != NULL_TREE) + { + bool type_err = false; + + if (TREE_CODE (TREE_TYPE (t)) != ARRAY_TYPE) + 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 (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/gfortran.h b/gcc/fortran/gfortran.h index 7bf1d5a0452..18e685ca1b1 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -1353,6 +1353,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; } @@ -1394,6 +1396,7 @@ enum OMP_LIST_NONTEMPORAL, OMP_LIST_ALLOCATE, OMP_LIST_HAS_DEVICE_ADDR, + OMP_LIST_USES_ALLOCATORS, OMP_LIST_NUM /* Must be the last. */ }; diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc index 714148138c2..64066047a33 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -948,6 +948,7 @@ enum omp_mask2 OMP_CLAUSE_ATTACH, OMP_CLAUSE_NOHOST, OMP_CLAUSE_HAS_DEVICE_ADDR, /* OpenMP 5.1 */ + OMP_CLAUSE_USES_ALLOCATORS, /* OpenMP 5.2 */ /* This must come last. */ OMP_MASK2_LAST }; @@ -1364,6 +1365,238 @@ gfc_match_omp_clause_reduction (char pc, gfc_omp_clauses *c, bool openacc, return MATCH_YES; } +/* uses_allocators ( allocator-list ) + + allocator-list: + allocator + allocator , allocator-list + allocator ( traits-array ) + allocator ( traits-array ) , allocator-list + + 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) +{ + gfc_symbol *sym; + gfc_symbol *memspace_sym= NULL; + gfc_symbol *traits_sym= NULL; + bool memspace_seen = false, traits_seen = false; + match m; + int i = 0; + + if (gfc_match ("uses_allocators ( ") != MATCH_YES) + return MATCH_NO; + + gfc_symbol *allocator_handle_kind, * memspace_handle_kind; + + if (gfc_get_symbol ("omp_allocator_handle_kind", NULL, &sym) + || !sym->value + || 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; + + if (gfc_get_symbol ("omp_memspace_handle_kind", NULL, &sym) + || !sym->value + || sym->value->expr_type != EXPR_CONSTANT + || sym->value->ts.type != BT_INTEGER) + { + gfc_error ("OpenMP % constant not found by " + "% clause at %C"); + goto error; + } + memspace_handle_kind = sym; + + 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 legacy " + "array 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 + || traits_sym->value == NULL + || !gfc_is_constant_expr (traits_sym->value)) + { + 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) + { + const char *pos; + 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. */ + || (pos = strstr (memspace_sym->name, "omp_")) == NULL + || pos != memspace_sym->name + || (pos = strstr (memspace_sym->name, "_mem_space")) == NULL + || *(pos + strlen ("_mem_space")) != '\0') + { + 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) + { + const char *pos; + /* Check if identifier is a 'omp_..._mem_alloc' pre-defined + allocator. */ + if ((pos = strstr (allocator_sym->name, "omp_")) == NULL + || pos != allocator_sym->name + || (pos = strstr (allocator_sym->name, "_mem_alloc")) == NULL + || *(pos + strlen ("_mem_alloc")) != '\0') + { + 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; + + return MATCH_YES; + + error: + return MATCH_ERROR; +} /* Match with duplicate check. Matches 'name'. If expr != NULL, it then matches '(expr)', otherwise, if open_parens is true, @@ -2924,6 +3157,9 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, ("use_device_addr (", &c->lists[OMP_LIST_USE_DEVICE_ADDR], false, NULL, NULL, true) == 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 @@ -3650,7 +3886,7 @@ cleanup: | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULTMAP \ | OMP_CLAUSE_IS_DEVICE_PTR | OMP_CLAUSE_IN_REDUCTION \ | OMP_CLAUSE_THREAD_LIMIT | OMP_CLAUSE_ALLOCATE \ - | OMP_CLAUSE_HAS_DEVICE_ADDR) + | OMP_CLAUSE_HAS_DEVICE_ADDR | 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) @@ -6282,7 +6518,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", "HAS_DEVICE_ADDR" }; + "NONTEMPORAL", "ALLOCATE", "HAS_DEVICE_ADDR", "USES_ALLOCATORS" }; STATIC_ASSERT (ARRAY_SIZE (clause_names) == OMP_LIST_NUM); if (omp_clauses == NULL) diff --git a/gcc/fortran/trans-array.cc b/gcc/fortran/trans-array.cc index 05134952db4..a2a2b889d03 100644 --- a/gcc/fortran/trans-array.cc +++ b/gcc/fortran/trans-array.cc @@ -6343,10 +6343,8 @@ 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; + + index = gfc_conv_mpz_to_tree (c->offset, gfc_index_integer_kind); if (mpz_cmp_si (c->repeat, 1) > 0) { diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc index 43d59abe9e0..1a3b1d5c347 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -2686,9 +2686,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); @@ -3657,6 +3664,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; } diff --git a/gcc/fortran/types.def b/gcc/fortran/types.def index cd79ad45167..18a1bec8724 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) @@ -154,6 +155,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.cc b/gcc/gimplify.cc index 2588824dce2..48a070443ba 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -9148,6 +9148,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, hash_map *struct_map_to_clause = NULL; hash_map *struct_seen_clause = NULL; hash_set *struct_deref_set = NULL; + + hash_set *allocate_clauses = NULL; + hash_set *uses_allocators_allocators = NULL; + tree *prev_list_p = NULL, *orig_list_p = list_p; int handled_depend_iterators = -1; int nowait = -1; @@ -9185,6 +9189,13 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, || code == OMP_TARGET_EXIT_DATA) omp_target_reorder_clauses (list_p); + if (code == OMP_TARGET + && (omp_requires_mask & OMP_REQUIRES_DYNAMIC_ALLOCATORS) == 0) + { + allocate_clauses = new hash_set (); + uses_allocators_allocators = new hash_set (); + } + while ((c = *list_p) != NULL) { bool remove = false; @@ -10884,6 +10895,18 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, OMP_CLAUSE_ALLOCATE_ALLOCATOR (c) = get_initialized_tmp_var (OMP_CLAUSE_ALLOCATE_ALLOCATOR (c), pre_p, NULL, false); + if (allocate_clauses + && OMP_CLAUSE_ALLOCATE_ALLOCATOR (c) + && DECL_P (OMP_CLAUSE_ALLOCATE_ALLOCATOR (c)) + && !allocate_clauses->contains (c)) + allocate_clauses->add (c); + break; + + case OMP_CLAUSE_USES_ALLOCATORS: + decl = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c); + if (uses_allocators_allocators + && !uses_allocators_allocators->contains (decl)) + uses_allocators_allocators->add (decl); break; case OMP_CLAUSE_DEFAULT: @@ -10936,6 +10959,26 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, list_p = &OMP_CLAUSE_CHAIN (c); } + if (code == OMP_TARGET + && (omp_requires_mask & OMP_REQUIRES_DYNAMIC_ALLOCATORS) == 0) + { + for (hash_set::iterator i = allocate_clauses->begin (); + i != allocate_clauses->end (); ++i) + { + tree allocator = OMP_CLAUSE_ALLOCATE_ALLOCATOR (*i); + if (uses_allocators_allocators->contains (allocator)) + continue; + + error_at (OMP_CLAUSE_LOCATION (*i), + "allocator %<%qE%>in % clause on target region " + "is missing % clause", + DECL_NAME (allocator), DECL_NAME (allocator)); + } + + delete allocate_clauses; + delete uses_allocators_allocators; + } + ctx->clauses = *orig_list_p; gimplify_omp_ctxp = ctx; if (struct_seen_clause) @@ -14165,6 +14208,79 @@ 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); + + /* + gimplify_assign (memspace_var, + fold_convert (pointer_sized_int_node, + memspace), + &init_seq); + */ + 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 cfa6483c7ae..e3103cea1c3 100644 --- a/gcc/omp-builtins.def +++ b/gcc/omp-builtins.def @@ -76,6 +76,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/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..3288181b1ec --- /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 uses_allocators (memspace(omp_high_bw_mem_space), traits(foo_traits) : bar) + ; + #pragma omp target 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\\(foo\\) uses_allocators\\(foo: memspace\\(\\), traits\\(\\)\\) private\\(bar\\) uses_allocators\\(bar: memspace\\(omp_high_bw_mem_space\\), traits\\(foo_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..74f411360c1 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/uses_allocators-2.c @@ -0,0 +1,32 @@ +/* { 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." } */ + ; + #pragma omp target uses_allocators (foo (xyz)) /* { dg-error "'xyz' undeclared .first use in this function." } */ + ; + #pragma omp target uses_allocators (foo (traits_array), baz (traits_array)) + ; + #pragma omp target uses_allocators (memspace(omp_no_such_space) : foo) /* { dg-error "'omp_no_such_space' undeclared" } */ + ; + #pragma omp target uses_allocators (memspace(1) : foo) /* { dg-error "expected '\\\)' before numeric constant" } */ + ; + #pragma omp target uses_allocators (memspace(omp_no_such_space) : foo, bar) /* { dg-error "memspace modifier must be constant enum of 'omp_memspace_handle_t' type" } */ + ; + #pragma omp target uses_allocators (traits(xyz) : bar) + ; + #pragma omp target uses_allocators (memspace(omp_high_bw_mem_space), traits(traits_array), memspace (omp_no_such_space) : bar) /* { dg-error "expected ':' before ',' token" } */ + ; + #pragma omp target uses_allocators (traitz(traits_array), memspace(omp_high_bw_mem_space) : bar) uses_allocators(foo) /* { dg-error "'traitz' undeclared .first use in this function." } */ + ; + return 0; +} 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..8afe9c4c331 --- /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 uses_allocators(memspace (omp_low_lat_mem_space) : bar) + block + end block + + !$omp target 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..530d604902f --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-2.f90 @@ -0,0 +1,44 @@ +! { 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 "'omp_non_existant_alloc' at .1. must be integer of 'omp_allocator_handle_kind' kind" } + block + end block + + !$omp target uses_allocators(bar(traits_array), foo (traits_array), ) ! { dg-error "Expected name of allocator at .1." } + block + end block + + !$omp target uses_allocators(traits(xyz) : bar) ! { dg-error "'xyz' at .1. must be of constant 'type.omp_alloctrait.' array type and have a constant initializer" } + block + end block + + !$omp target uses_allocators(memspace(omp_non_existant_mem_space) : foo) ! { dg-error "'omp_non_existant_mem_space' at .1. is not a pre-defined memory space name" } + block + end block + + !$omp target uses_allocators(traits(traits_array), traits(traits_array) : bar) ! { dg-error "Multiple traits modifiers at .1." } + block + end block + + !$omp target uses_allocators(memspace(omp_default_mem_space), memspace(omp_default_mem_space) : foo) ! { dg-error "Multiple memspace modifiers at .1." } + block + end block + + !$omp target uses_allocators(memspace(omp_default_mem_space), traits(traits_array), traits(traits_array) : foo) ! { dg-error "Only two modifiers are allowed on 'uses_allocators' clause at .1." } + block + end block + + !$omp target uses_allocators(memspace(omp_default_mem_space), traits(traits_array) : foo, bar) ! { dg-error "When using modifiers, only a single allocator can be specified in each 'uses_allocators' clause 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..064ccf455b1 --- /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''in 'allocate' clause on target region is missing 'uses_allocators.bar.' clause" } + block + allocate(arr(100)) + end block + +end program main diff --git a/gcc/tree-core.h b/gcc/tree-core.h index f1c2b6413a3..7ac0b47ac2d 100644 --- a/gcc/tree-core.h +++ b/gcc/tree-core.h @@ -522,6 +522,9 @@ enum omp_clause_code { /* OpenACC clause: nohost. */ OMP_CLAUSE_NOHOST, + + /* OpenMP clause: uses_allocators. */ + OMP_CLAUSE_USES_ALLOCATORS, }; #undef DEFTREESTRUCT diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc index 99af977979d..a46db024157 100644 --- a/gcc/tree-pretty-print.cc +++ b/gcc/tree-pretty-print.cc @@ -769,6 +769,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.cc b/gcc/tree.cc index 4cf3785270b..973a8366372 100644 --- a/gcc/tree.cc +++ b/gcc/tree.cc @@ -349,6 +349,7 @@ unsigned const char omp_clause_num_ops[] = 0, /* OMP_CLAUSE_IF_PRESENT */ 0, /* OMP_CLAUSE_FINALIZE */ 0, /* OMP_CLAUSE_NOHOST */ + 3, /* OMP_CLAUSE_USES_ALLOCATORS */ }; const char * const omp_clause_code_name[] = @@ -439,6 +440,7 @@ const char * const omp_clause_code_name[] = "if_present", "finalize", "nohost", + "uses_allocators", }; /* Unless specific to OpenACC, we tend to internally maintain OpenMP-centric diff --git a/gcc/tree.h b/gcc/tree.h index 8844471e9a5..bfe2cd82232 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -1824,6 +1824,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)