From patchwork Tue May 10 11:29:23 2022 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Chung-Lin Tang X-Patchwork-Id: 53734 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 831B3395A00C for ; Tue, 10 May 2022 11:29:56 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id B94BB3955CBC for ; Tue, 10 May 2022 11:29:32 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org B94BB3955CBC 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,214,1647331200"; d="scan'208";a="75521953" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa3.mentor.iphmx.com with ESMTP; 10 May 2022 03:29:32 -0800 IronPort-SDR: o47Y4fxbHcB6PQzHy5kU1s55uj2iC7i8Nwc0hQFCal5+epyCZ/srHYsAUi7F1acRSF2fbgHtwn N5c1XBF3HgV+ikfixKufy0lIz5ZHSYAWc79F6npRX9L5QaZLdsjzNno7qcfQIZJ4SMmfOraYNB 5WlFENBdEH1Hnt9CxhJmN465pGI9/Vcpb5jqz1tSdHBYUf6u85lBceokrol5nzCLtOV6UnmFNT ZQv1RxgF3+4SZmU7is/EXTsFqZydDwqtQ3YxFlwnW9i9rvQ/OFe7boMNMlOhGdXLc+evXCuQ8d VVY= Message-ID: Date: Tue, 10 May 2022 19:29:23 +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 Subject: [PATCH, OpenMP, v2] Implement uses_allocators clause for target regions Content-Language: en-US To: Tobias Burnus , gcc-patches , Jakub Jelinek , Catherine Moore , Andrew Stubbs , Hafiz Abid Qadeer References: <46d77e14-080c-db6c-4032-e12899c5d059@codesourcery.com> <9c0945fa-1054-095e-86ae-a9d8dd1ab625@codesourcery.com> From: Chung-Lin Tang In-Reply-To: <9c0945fa-1054-095e-86ae-a9d8dd1ab625@codesourcery.com> X-ClientProxiedBy: svr-orw-mbx-14.mgc.mentorg.com (147.34.90.214) To svr-orw-mbx-02.mgc.mentorg.com (147.34.90.202) X-Spam-Status: No, score=-8.9 required=5.0 tests=BAYES_00, BODY_8BITS, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, KAM_MANYTO, 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" On 2022/5/7 12:40 AM, Tobias Burnus wrote: > > Can please also handle the new clause in Fortran's dump-parse-tree.cc? > > I did see some split handling in C, but not in Fortran; do you also need > to up update gfc_split_omp_clauses in Fortran's trans-openmp.cc? Done. > Actually, glancing at the testcases, no combined construct (like > "omp target parallel") is used, I think that would be useful because of ↑. Okay, added some to testcases. >> +/* OpenMP 5.2: >> +   uses_allocators ( allocator-list ) > That's not completely true: uses_allocators is OpenMP 5.1. > However, 5.1 only supports (for non-predefined allocators): >    uses_allocators( allocator(traits) ) > while OpenMP 5.2 added modifiers: >    uses_allocatrors( traits(...), memspace(...) : allocator ) > and deprecated the 5.1 'allocator(traits)'. (Scheduled for removal in OMP 6.0) > > The advantage of 5.2 syntax is that a memory space can be defined. I supported both syntaxes, that's why I designated it as "5.2". > BTW: This makes uses_allocators the first OpenMP 5.2 feature which > will make it into GCC :-) :) > > gcc/fortran/openmp.cc: >> +  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; > I think you rather want to use >   gfc_find_symbol ("omp_...", NULL, true, &sym) >   || sym == NULL > where true is for parent_flag to search also the parent namespace. > (The function returns 1 if the symbol is ambiguous, 0 otherwise - > including 0 + sym == NULL when the symbol could not be found.) > >   || sym->attr.flavor != FL_PARAMETER >   || sym->ts.type != BT_INTEGER >   || sym->attr.dimension > > Looks cleaner than to access sym->value. The attr.dimension is just > to makes sure the user did not smuggle an array into this. > (Invalid as omp_... is a reserved namespace but users will still do > this and some are good in finding ICE as hobby.) Well, the intention here is to search for "omp_allocator_handle_kind" and "omp_memspace_handle_kind", and use their value to check if the kinds are the same as declared allocator handles and memspace constant. Not to generally search for "omp_...". However the sym->attr.dimension test seems useful, added in new v2 patch. > However, I fear that will fail for the following two examples (both untested): > >   use omp_lib, my_kind = omp_allocator_handle_kind >   integer(my_kind) :: my_allocator > > as this gives 'my_kind' in the symtree->name (while symtree->n.sym->name is "omp_..."). > Hence, by searching the symtree for 'omp_...' the symbol will not be found. > > > It will likely also fail for the following more realistic example: ... > subroutine foo >   use m >   use omp_lib, only: omp_alloctrait ... >   !$omp target uses_allocators(my_allocator(traits_array) allocate(my_allocator:A) firstprivate(A) >      ... >   !$omp end target > end If someone wants to use OpenMP allocators, but intentionally only imports insufficient standard symbols from omp_lib, then he/she is on their own :) The specification really makes this quite clear: omp_allocator_handle_kind, omp_alloctrait, omp_memspace_handle_kind are all part of the same package. > In this case, omp_allocator_handle_kind is not in the namespace of 'foo' > but the code should be still valid. Thus, an alternative would be to hard-code > the value - as done for the depobj. As we have: > >         integer, parameter :: omp_allocator_handle_kind = c_intptr_t >         integer, parameter :: omp_memspace_handle_kind = c_intptr_t > > that would be >    sym->ts.type == BT_CHARACTER >    sym->ts.kind == gfc_index_integer_kind > for the allocator variable and the the memspace kind. > > However, I grant that either example is not very typical. The second one is more > natural – such a code will very likely be written in the real world. But not > with uses_allocators but rather with "!$omp requires dynamic_allocators" and > omp_init_allocator(). > > Thoughts? As above. I mean, what is so hard with including "use omp_lib" where you need it? :D > * * * > > gcc/fortran/openmp.cc >> +      if (++i > 2) >> +    { >> +      gfc_error ("Only two modifiers are allowed on % " >> +             "clause at %C"); >> +      goto error; >> +    } >> + > > Is this really needed? There is a check for multiple traits and multiple memspace > Thus, 'trait(),memspace(),trait()' is already handled and > 'trait(),something' give a break and will lead to an error as in that case > a ':' and not ',something' is expected. I think it could be worth reminding that limitation, instead of a generic error. >> +      if (gfc_match_char ('(') == MATCH_YES) >> +    { >> +      if (memspace_seen || traits_seen) >> +        { >> +          gfc_error ("Modifiers cannot be used with legacy " >> +             "array syntax at %C"); > I wouldn't uses the term 'array synax' to denote >   uses_allocators(allocator (alloc_array) ) > How about: >   error: "Using both modifiers and allocator variable with traits argument" > > (And I think 'deprecated' is better than 'legacy', if we really want to use it.) I've changed it to "(deprecated) traits array list syntax", is that better? >> +      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)) > > I think the gfc_is_constant_expr is unreachable as you already > have checked FL_PARAMETER. Thus, you can remove the last two > lines. Okay. > [Regarding the traits_sym->ts.u.derived->name, I am not sure whether that > won't fail with >   use omp_lib, trait_t => omp_alloctrait > but I have not checked. It likely does work correctly.] > >> +          /* 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') > > I wonder whether that's not more readable written as: >    || !startswith (memspace_sym->name, "omp_") >    || !endswith (memspace_sym->name, "_mem_space") Thanks, didn't know it was this convenient :) I've attached v2 of the patch. Currently in testing. Thanks, Chung-Lin 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..bbdec92780b 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,233 @@ 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 t = NULL_TREE, nl = list; + 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); + 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); + } + + if (!parens2.require_close (parser)) + { + parens.skip_until_found_close (parser); + return list; + } + + 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 (t == error_mark_node) + return list; + } + has_modifiers = true; + } + } + + 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 if (c_parser_next_token_is (parser, CPP_COMMA)) + c_parser_error (parser, "modifiers can only be used with " + "a single allocator in % " + "clause"); + } + else + c_parser_error (parser, "expected identifier"); + } + else + { + 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); + t = error_mark_node; + } + c_parser_consume_token (parser); + + traits_var = NULL_TREE; + 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); + traits_var = error_mark_node; + } + c_parser_consume_token (parser); + } + else + c_parser_error (parser, "expected identifier"); + parens2.require_close (parser); + } + + if (t != error_mark_node && traits_var != 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) = 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 +17279,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 +21294,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..e041bc669a9 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,247 @@ 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 = list; + 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) + { + 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 (!parens2.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; + } + } + has_modifiers = true; + } + } + + 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 if (cp_lexer_next_token_is (parser->lexer, CPP_COMMA)) + cp_parser_error (parser, "modifiers can only be used with " + "a single allocator in % " + "clause"); + } + else + cp_parser_error (parser, "expected identifier"); + } + else + { + 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); + + traits_var = NULL_TREE; + 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); + cp_lexer_consume_token (parser->lexer); + } + else + cp_parser_error (parser, "expected identifier"); + parens2.require_close (parser); + } + + if (t != error_mark_node && traits_var != 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) = 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); + } + } + + cp_parser_skip_to_closing_parenthesis (parser, /*recovering=*/false, + /*or_comma=*/false, + /*consume_paren=*/true); + return nl; +} + /* OpenMP 2.5: lastprivate ( variable-list ) @@ -40283,6 +40526,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 +44538,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..c4ff73e7899 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 %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.cc b/gcc/fortran/dump-parse-tree.cc index 3635460bffd..3ac7fc846ac 100644 --- a/gcc/fortran/dump-parse-tree.cc +++ b/gcc/fortran/dump-parse-tree.cc @@ -1423,6 +1423,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); @@ -1689,6 +1703,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 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..a187e75e1fe 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,234 @@ 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_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; + + 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; + } + 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 (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 (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; + + 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 +3153,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 +3882,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 +6514,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..b094b17f054 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; } @@ -6074,6 +6104,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 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..3e858fa9512 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,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 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..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..78a2d786248 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/uses_allocators-2.c @@ -0,0 +1,37 @@ +/* { 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 '\\\)' before numeric constant" } */ + ; /* { dg-error "expected '#pragma omp' clause before ':' token" "" { target *-*-* } .-1 } */ + #pragma omp target uses_allocators (memspace(omp_no_such_space) : foo, bar) /* { dg-error "modifiers can only be used with a single allocator in 'uses_allocators' clause" } */ + ; /* { dg-error "memspace modifier must be constant enum of 'omp_memspace_handle_t' type" "" { target c } .-1 } */ + /* { dg-error "'omp_no_such_space' has not been declared" "" { target c++ } .-2 } */ + #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 "expected ':' before ',' token" } */ + ; + #pragma omp target uses_allocators (traitz(traits_array), memspace(omp_high_bw_mem_space) : bar) /* { dg-error "'traitz' undeclared" "" { target c } } */ + ; /* { dg-error "'memspace' undeclared" "" { target c } .-1 } */ + /* { dg-error "expected '\\\)' before ':' token" "" { target c } .-2 } */ + /* { dg-error "'traitz' has not been declared" "" { target c++ } .-3 } */ + /* { dg-error "'memspace' has not been declared" "" { target c++ } .-4 } */ + 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..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..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)