From patchwork Mon Nov 20 10:42:02 2023 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Tobias Burnus X-Patchwork-Id: 167032 Return-Path: Delivered-To: ouuuleilei@gmail.com Received: by 2002:a59:9910:0:b0:403:3b70:6f57 with SMTP id i16csp2109531vqn; Mon, 20 Nov 2023 02:42:47 -0800 (PST) X-Google-Smtp-Source: AGHT+IFNxhvifjZVDPVO+ugBIGtyG2aMeqDOeDKaXFWIs6nJliA4tEL8JeezK4AidO/4j44r/v2x X-Received: by 2002:a05:6214:48f:b0:671:7955:606e with SMTP id pt15-20020a056214048f00b006717955606emr10626500qvb.49.1700476967386; Mon, 20 Nov 2023 02:42:47 -0800 (PST) ARC-Seal: i=2; a=rsa-sha256; t=1700476967; cv=pass; d=google.com; s=arc-20160816; b=ZQSyu1pijimLkwW8WosXFdCZiUgXaRW8k9r3GbRFpk6LxFjAZrI0VQNncBispBAgVt ZHXwJ7P5WNUYXYxj9LNP59OWxytnblTKWbaveHvMabIlGUxK+aBjK70B077pCWBn/UXH H2dD8yUYQkbtR7D0PttKcvQ/Ja4Wk+yUM4Pfdsv6ZCsHhsRN2PPEh/jW5RBCStKe1Rsg 9naRzxLjgSY7w0ULtRhlJecicUPtzMPrBL7jtQw2KsSpWhSEkuqUwnVBxUx3PKOcqWf3 tlihqb7dZgpRNH04WZmuhgy7wT2LMO8XXABiTOZ97TLFZ2jgisr3qjOfgHTjOY7YdXkc 8JXA== ARC-Message-Signature: i=2; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20160816; h=errors-to:list-subscribe:list-help:list-post:list-archive :list-unsubscribe:list-id:precedence:subject:from:to :content-language:user-agent:mime-version:date:message-id :ironport-sdr:arc-filter:dmarc-filter:delivered-to; bh=zhqZE4zdq6OUUFtyl5ETDTTss2SsJdOUu+NGrYWRpDE=; fh=ARF5Mvhgv4mOzEi1LJHgn82UOzar4LOruqM9inFDcVw=; b=QkMC1W6iKMtNoVO3gnLVZ+LPH8S6RVYttGY/GkZkfmVvI5ir5jjLGkVrlonpVuQv+M hb63XsXjUSDDDm6sDFwzcYOlgDDh7X6Vc3f4pYtph/UHbRVnSBoADBpQsLCqz+MWS2Iu FxhtJKJatTKMt+w9g/Ro2SiQgtCGkztm0F5sHdHvtUUTfHbtP94jlEIc2ulMNj6NhOGh mHzYpKsePMut9nZl5VpCjP3km3xqgYV7r/b00Nk4eqhu4ELbKmVKECW4TE9hs8W/dWYW GP4S+YCn/r1bcy/gu5ZulYcllhVlWLRF2Oit3ja62/ZXdJ1ZBEsXXxa15+8fd2togCDM FHFA== ARC-Authentication-Results: i=2; mx.google.com; arc=pass (i=1); spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 8.43.85.97 as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" Received: from server2.sourceware.org (server2.sourceware.org. [8.43.85.97]) by mx.google.com with ESMTPS id o3-20020a0cecc3000000b00679da1cab69si1818911qvq.520.2023.11.20.02.42.47 for (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Mon, 20 Nov 2023 02:42:47 -0800 (PST) Received-SPF: pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 8.43.85.97 as permitted sender) client-ip=8.43.85.97; Authentication-Results: mx.google.com; arc=pass (i=1); spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 8.43.85.97 as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 5DE2B3857BA7 for ; Mon, 20 Nov 2023 10:42:42 +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 B5D5F3858C98; Mon, 20 Nov 2023 10:42:08 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org B5D5F3858C98 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com ARC-Filter: OpenARC Filter v1.0.0 sourceware.org B5D5F3858C98 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=68.232.137.180 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1700476934; cv=none; b=DWyQuEY4ArTrjqwEKwd3rg2P2jq+5YyoYImvLlxgJQ7qFA5OzlFGGGqc0mKpVVUj4mw0zg9cecfldQVJN7AiZw7Mb8QzXToe42VPYUlLWJ5QoU1qTrGJF2HgG+duuNzSF+cdr8MGvYEqj0Zy/YNaDwJ8CgpkemiqLcg+hRO0JDU= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1700476934; c=relaxed/simple; bh=zhqZE4zdq6OUUFtyl5ETDTTss2SsJdOUu+NGrYWRpDE=; h=Message-ID:Date:MIME-Version:To:From:Subject; b=etN4uPFAA4w+rHvDL0Q2oonxdabk+CmX/I8E8gGODkGDXx6Aq2Kb+Ee63iJyMGYAZLDoBgBYkL2CJKiO78Tpsp0I820kqAc1dh9qWH6kVcTfaOmlnKGoqJJkEw5OBnFJVFcGfRYElASbTgHIZrTMOhZ9dLJJBLpGTvBGZm6voNs= ARC-Authentication-Results: i=1; server2.sourceware.org X-CSE-ConnectionGUID: Xo9Nu9cLSA+N5Ij1Qfwkvw== X-CSE-MsgGUID: avhWmsKaQKKH9CjyRt73gA== X-IronPort-AV: E=Sophos;i="6.04,213,1695715200"; d="diff'?scan'208";a="23243078" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa3.mentor.iphmx.com with ESMTP; 20 Nov 2023 02:42:07 -0800 IronPort-SDR: TZwrPJLPzM+sGpQoxWOO467d0Eea6qKGF3SdQnn+peTsU6vcCEAp9kgJfIami+gQBIzV3TvX7X gcOA7K6nDqFkhPoAfo2vbu74ZLtpW8dFfO2JbEEn5pXwpWHPTV5GBzC3ixIsf6iiNB06sjSnNh WkazWwWTaURiJ3cQSgoUd7mUSK8qwMndqb5Yluok2APY9VqgWvEXWtnUqrhawmgUVFFpaJwk3h I4osJtgRwA43XsSFOUzDH1spxiTpoDhwLnIKKI8SrzDm+wsB2oTsDPBwmunWHqwNir/hppM1Z2 B3E= Message-ID: <545c0550-b8ed-41f4-bfd7-4e79aaf6be79@codesourcery.com> Date: Mon, 20 Nov 2023 11:42:02 +0100 MIME-Version: 1.0 User-Agent: Mozilla Thunderbird Content-Language: en-US To: gcc-patches , fortran From: Tobias Burnus Subject: [patch] OpenMP: Add uses_allocators support X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-10.mgc.mentorg.com (139.181.222.10) To svr-ies-mbx-12.mgc.mentorg.com (139.181.222.12) X-Spam-Status: No, score=-11.3 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, KAM_SHORT, SPF_HELO_PASS, SPF_PASS, TXREP, T_SCC_BODY_TEXT_LINE autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org X-getmail-retrieved-from-mailbox: INBOX X-GMAIL-THRID: 1783079336251876189 X-GMAIL-MSGID: 1783079336251876189 This adds middle end support for uses_allocators, wires Fortran to use it and add C/C++ parsing support. This is based on Chung-Lin's patch at https://gcc.gnu.org/pipermail/gcc-patches/2022-June/596412.html There is a known firstprivate/private issue for C/C++, see: https://gcc.gnu.org/PR110347 but that one is unrelated. Tested on x86-64-gnu-linux w/o offloading; with offloading still running. Comments, suggestions? Tobias ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 OpenMP: Add uses_allocators support 2023-11-19 Tobias Burnus Chung-Lin Tang gcc/ChangeLog: * builtin-types.def (BT_FN_VOID_PTRMODE): (BT_FN_PTRMODE_PTRMODE_INT_PTR): Add. * gimplify.cc (gimplify_bind_expr): Diagnose missing uses_allocators clause. (gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses, gimplify_omp_workshare): Handle uses_allocators. * omp-builtins.def (BUILT_IN_OMP_INIT_ALLOCATOR, BUILT_IN_OMP_DESTROY_ALLOCATOR): Add. * omp-low.cc (scan_sharing_clauses): * tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_USES_ALLOCATORS. * tree.cc (omp_clause_num_ops, omp_clause_code_name): Likewise. * tree-pretty-print.cc (dump_omp_clause): Handle it. * tree.h (OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR, OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE, OMP_CLAUSE_USES_ALLOCATORS_TRAITS): New. gcc/c-family/ChangeLog: * c-omp.cc (c_omp_split_clauses): Hande uses_allocators. * c-pragma.h (enum pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_USES_ALLOCATORS. gcc/c/ChangeLog: * c-parser.cc (c_parser_omp_clause_uses_allocators): New. (c_parser_omp_clause_name, c_parser_omp_all_clauses, OMP_TARGET_CLAUSE_MASK): Handle uses_allocators. * c-typeck.cc (c_finish_omp_clauses): Likewise. gcc/cp/ChangeLog: * parser.cc (cp_parser_omp_clause_uses_allocators): New. (cp_parser_omp_clause_name, cp_parser_omp_all_clauses, OMP_TARGET_CLAUSE_MASK): Handle uses_allocators. * semantics.cc (finish_omp_clauses): Likewise. gcc/fortran/ChangeLog: * trans-array.cc (gfc_conv_array_initializer): Set PURPOSE when building constructor for get_initialized_tmp_var. * trans-openmp.cc (gfc_trans_omp_clauses): Handle uses_allocators. * types.def (BT_FN_VOID_PTRMODE, BT_FN_PTRMODE_PTRMODE_INT_PTR): Add. libgomp/ChangeLog: * testsuite/libgomp.c++/c++.exp (check_effective_target_c, check_effective_target_c++): Add. * testsuite/libgomp.c/c.exp (check_effective_target_c, check_effective_target_c++): Add. * testsuite/libgomp.fortran/uses_allocators_2.f90: Remove 'sorry'. * testsuite/libgomp.c-c++-common/uses_allocators-1.c: New test. * testsuite/libgomp.c-c++-common/uses_allocators-2.c: New test. * testsuite/libgomp.c-c++-common/uses_allocators-3.c: New test. * testsuite/libgomp.c-c++-common/uses_allocators-4.c: New test. * testsuite/libgomp.fortran/uses_allocators_3.f90: New test. * testsuite/libgomp.fortran/uses_allocators_4.f90: New test. * testsuite/libgomp.fortran/uses_allocators_5.f90: New test. * testsuite/libgomp.fortran/uses_allocators_6.f90: New test. gcc/testsuite/ChangeLog: * gfortran.dg/gomp/allocate-1.f90: Add uses_allocators. * gfortran.dg/gomp/scope-6.f90: Update dg-scan-tree-dump. * 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. Co-authored-by: Chung-Lin Tang gcc/builtin-types.def | 3 + gcc/c-family/c-omp.cc | 1 + gcc/c-family/c-pragma.h | 1 + gcc/c/c-parser.cc | 216 ++++++++++++++++++- gcc/c/c-typeck.cc | 105 +++++++++ gcc/cp/parser.cc | 237 ++++++++++++++++++++- gcc/cp/semantics.cc | 95 +++++++++ gcc/fortran/trans-array.cc | 5 +- gcc/fortran/trans-openmp.cc | 42 +++- gcc/fortran/types.def | 3 + gcc/gimplify.cc | 183 +++++++++++++++- gcc/omp-builtins.def | 4 + gcc/omp-low.cc | 32 +++ .../c-c++-common/gomp/uses_allocators-1.c | 46 ++++ .../c-c++-common/gomp/uses_allocators-2.c | 33 +++ gcc/testsuite/gfortran.dg/gomp/allocate-1.f90 | 7 +- gcc/testsuite/gfortran.dg/gomp/scope-6.f90 | 2 +- .../gfortran.dg/gomp/uses_allocators-1.f90 | 23 ++ gcc/tree-core.h | 3 + gcc/tree-pretty-print.cc | 14 ++ gcc/tree.cc | 2 + gcc/tree.h | 9 + libgomp/testsuite/libgomp.c++/c++.exp | 9 + .../libgomp.c-c++-common/uses_allocators-1.c | 53 +++++ .../libgomp.c-c++-common/uses_allocators-2.c | 39 ++++ .../libgomp.c-c++-common/uses_allocators-3.c | 37 ++++ .../libgomp.c-c++-common/uses_allocators-4.c | 53 +++++ libgomp/testsuite/libgomp.c/c.exp | 8 + .../libgomp.fortran/uses_allocators_2.f90 | 22 +- .../libgomp.fortran/uses_allocators_3.f90 | 62 ++++++ .../libgomp.fortran/uses_allocators_4.f90 | 54 +++++ .../libgomp.fortran/uses_allocators_5.f90 | 14 ++ .../libgomp.fortran/uses_allocators_6.f90 | 50 +++++ 33 files changed, 1435 insertions(+), 32 deletions(-) diff --git a/gcc/builtin-types.def b/gcc/builtin-types.def index 43381bc8949..da9719272ec 100644 --- a/gcc/builtin-types.def +++ b/gcc/builtin-types.def @@ -385,6 +385,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) @@ -829,6 +830,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 95b6c1e623f..82422fc8c3e 100644 --- a/gcc/c-family/c-omp.cc +++ b/gcc/c-family/c-omp.cc @@ -2026,6 +2026,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_DOACROSS: diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h index 98177913053..7b1d9dc7232 100644 --- a/gcc/c-family/c-pragma.h +++ b/gcc/c-family/c-pragma.h @@ -161,6 +161,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 703f9570dbc..db69a418f59 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -14722,6 +14722,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)) @@ -17568,6 +17570,213 @@ c_parser_omp_clause_allocate (c_parser *parser, tree list) return nl; } +/* OpenMP 5.0: + uses_allocators ( allocator-list ) + + allocator-list: + allocator + allocator , allocator-list + allocator ( traits-array ) + allocator ( traits-array ) , allocator-list + + OpenMP 5.2: + + uses_allocators ( modifier : allocator-list ) + uses_allocators ( modifier , modifier : allocator-list ) + + modifier: + traits ( traits-array ) + memspace ( mem-space-handle ) */ + +static tree +c_parser_omp_clause_uses_allocators (c_parser *parser, tree list) +{ + location_t clause_loc = c_parser_peek_token (parser)->location; + tree t = NULL_TREE, nl = list; + matching_parens parens; + if (!parens.require_open (parser)) + return list; + + tree memspace_expr = NULL_TREE; + tree traits_var = NULL_TREE; + + struct item_tok + { + location_t loc; + tree id; + item_tok (void) : loc (UNKNOWN_LOCATION), id (NULL_TREE) {} + }; + struct item { item_tok name, arg; }; + auto_vec *modifiers = NULL, *allocators = NULL; + auto_vec *cur_list = new auto_vec (4); + + while (true) + { + item it; + + if (c_parser_next_token_is (parser, CPP_NAME)) + { + c_token *tok = c_parser_peek_token (parser); + it.name.id = tok->value; + it.name.loc = tok->location; + c_parser_consume_token (parser); + + if (c_parser_next_token_is (parser, CPP_OPEN_PAREN)) + { + matching_parens parens2; + parens2.consume_open (parser); + + if (c_parser_next_token_is (parser, CPP_NAME)) + { + tok = c_parser_peek_token (parser); + it.arg.id = tok->value; + it.arg.loc = tok->location; + c_parser_consume_token (parser); + } + else + { + c_parser_error (parser, "expected identifier"); + parens2.skip_until_found_close (parser); + goto end; + } + parens2.skip_until_found_close (parser); + } + } + + cur_list->safe_push (it); + + if (c_parser_next_token_is (parser, CPP_COMMA)) + c_parser_consume_token (parser); + else if (c_parser_next_token_is (parser, CPP_COLON)) + { + if (modifiers) + { + c_parser_error (parser, "expected %<)%>"); + goto end; + } + else + { + c_parser_consume_token (parser); + modifiers = cur_list; + cur_list = new auto_vec (4); + } + } + else if (c_parser_next_token_is (parser, CPP_CLOSE_PAREN)) + { + gcc_assert (allocators == NULL); + allocators = cur_list; + cur_list = NULL; + break; + } + else + { + c_parser_error (parser, "expected %<)%>"); + goto end; + } + } + + if (modifiers) + for (unsigned i = 0; i < modifiers->length (); i++) + { + item& it = (*modifiers)[i]; + const char *p = IDENTIFIER_POINTER (it.name.id); + int strcmp_traits = 1, strcmp_memspace = 1; + + if ((strcmp_traits = strcmp ("traits", p)) == 0 + || (strcmp_memspace = strcmp ("memspace", p)) == 0) + { + if ((strcmp_traits == 0 && traits_var != NULL_TREE) + || (strcmp_memspace == 0 && memspace_expr != NULL_TREE)) + { + error_at (it.name.loc, "duplicate %qs modifier", p); + goto end; + } + t = lookup_name (it.arg.id); + if (t == NULL_TREE) + { + undeclared_variable (it.arg.loc, it.arg.id); + t = error_mark_node; + } + else if (strcmp_memspace == 0) + memspace_expr = t; + else if (strcmp_traits == 0) + traits_var = t; + else + gcc_unreachable (); + } + else + { + error_at (it.name.loc, "unknown modifier %qE", it.name.id); + goto end; + } + } + + if (allocators) + { + if (modifiers) + { + if (allocators->length () > 1) + { + error_at ((*allocators)[1].name.loc, + "% clause only accepts a single " + "allocator when using modifiers"); + goto end; + } + else if ((*allocators)[0].arg.id) + { + error_at ((*allocators)[0].arg.loc, + "legacy %<%E(%E)%> traits syntax not allowed in " + "% clause when using modifiers", + (*allocators)[0].name.id, (*allocators)[0].arg.id); + goto end; + } + } + + for (unsigned i = 0; i < allocators->length (); i++) + { + item& it = (*allocators)[i]; + t = lookup_name (it.name.id); + if (t == NULL_TREE) + { + undeclared_variable (it.name.loc, it.name.id); + goto end; + } + else if (t != error_mark_node) + { + tree t2 = NULL_TREE; + if (it.arg.id) + { + t2 = lookup_name (it.arg.id); + if (t2 == NULL_TREE) + { + undeclared_variable (it.arg.loc, it.arg.id); + goto end; + } + } + else + t2 = traits_var; + + tree c = build_omp_clause (clause_loc, + OMP_CLAUSE_USES_ALLOCATORS); + OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c) = t; + OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c) = memspace_expr; + OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c) = t2; + OMP_CLAUSE_CHAIN (c) = nl; + nl = c; + } + } + } + end: + if (cur_list) + delete cur_list; + if (modifiers) + delete modifiers; + if (allocators) + delete allocators; + parens.skip_until_found_close (parser); + return nl; +} + /* OpenMP 4.0: linear ( variable-list ) linear ( variable-list : expression ) @@ -19237,6 +19446,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"; @@ -23648,7 +23861,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 1dbb4471a88..e17414a5245 100644 --- a/gcc/c/c-typeck.cc +++ b/gcc/c/c-typeck.cc @@ -15136,6 +15136,111 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) break; } gcc_unreachable (); + + case OMP_CLAUSE_USES_ALLOCATORS: + t = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c); + if ((VAR_P (t) || TREE_CODE (t) == PARM_DECL) + && (bitmap_bit_p (&generic_head, DECL_UID (t)) + || bitmap_bit_p (&map_head, DECL_UID (t)) + || bitmap_bit_p (&firstprivate_head, DECL_UID (t)) + || bitmap_bit_p (&lastprivate_head, DECL_UID (t)))) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qE appears more than once in data clauses", t); + remove = true; + break; + } + else + bitmap_set_bit (&generic_head, DECL_UID (t)); + if (TREE_CODE (TREE_TYPE (t)) != ENUMERAL_TYPE + || strcmp (IDENTIFIER_POINTER (TYPE_IDENTIFIER (TREE_TYPE (t))), + "omp_allocator_handle_t") != 0) + { + error_at (OMP_CLAUSE_LOCATION (c), + "allocator must be of % type"); + remove = true; + break; + } + if (TREE_CODE (t) == CONST_DECL) + { + /* Currently for pre-defined allocators in libgomp, we do not + require additional init/fini inside target regions, so discard + such clauses. */ + remove = true; + + if (strcmp (IDENTIFIER_POINTER (DECL_NAME (t)), + "omp_null_allocator") == 0) + { + error_at (OMP_CLAUSE_LOCATION (c), + "% cannot be used in " + "% clause"); + break; + } + + if (OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c) + || OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c)) + { + error_at (OMP_CLAUSE_LOCATION (c), + "modifiers cannot be used with pre-defined " + "allocators"); + break; + } + } + t = OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c); + if (t != NULL_TREE + && (TREE_CODE (t) != CONST_DECL + || TREE_CODE (TREE_TYPE (t)) != ENUMERAL_TYPE + || strcmp (IDENTIFIER_POINTER (TYPE_IDENTIFIER (TREE_TYPE (t))), + "omp_memspace_handle_t") != 0)) + { + error_at (OMP_CLAUSE_LOCATION (c), "memspace modifier must be " + "constant enum of % type"); + remove = true; + break; + } + t = OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c); + if (t != NULL_TREE) + { + bool type_err = false; + + if (TREE_CODE (TREE_TYPE (t)) != ARRAY_TYPE + || DECL_SIZE (t) == NULL_TREE) + type_err = true; + else + { + tree elem_t = TREE_TYPE (TREE_TYPE (t)); + if (TREE_CODE (elem_t) != RECORD_TYPE + || strcmp (IDENTIFIER_POINTER (TYPE_IDENTIFIER (elem_t)), + "omp_alloctrait_t") != 0 + || !TYPE_READONLY (elem_t)) + type_err = true; + } + if (type_err) + { + if (TREE_CODE (t) != ERROR_MARK) + error_at (OMP_CLAUSE_LOCATION (c), "traits array %qE must " + "be of % type", t); + else + 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; + pc = &OMP_CLAUSE_CHAIN (c); + continue; case OMP_CLAUSE_DEPEND: case OMP_CLAUSE_AFFINITY: t = OMP_CLAUSE_DECL (c); diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index f556b8f3c01..15216a08b54 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -37657,6 +37657,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)) @@ -39983,6 +39985,234 @@ cp_parser_omp_clause_allocate (cp_parser *parser, tree list) return nlist; } +/* OpenMP 5.0: + uses_allocators ( allocator-list ) + + allocator-list: + allocator + allocator , allocator-list + allocator ( traits-array ) + allocator ( traits-array ) , allocator-list + + OpenMP 5.2: + + uses_allocators ( modifier : allocator-list ) + uses_allocators ( modifier , modifier : allocator-list ) + + modifier: + traits ( traits-array ) + memspace ( mem-space-handle ) */ + +static tree +cp_parser_omp_clause_uses_allocators (cp_parser *parser, tree list) +{ + location_t clause_loc + = cp_lexer_peek_token (parser->lexer)->location; + tree t = NULL_TREE, nl = list; + matching_parens parens; + if (!parens.require_open (parser)) + return list; + + tree memspace_expr = NULL_TREE; + tree traits_var = NULL_TREE; + + struct item_tok + { + location_t loc; + tree id; + item_tok (void) : loc (UNKNOWN_LOCATION), id (NULL_TREE) {} + }; + struct item { item_tok name, arg; }; + auto_vec *modifiers = NULL, *allocators = NULL; + auto_vec *cur_list = new auto_vec (4); + + while (true) + { + item it; + + if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)) + { + cp_token *tok = cp_lexer_peek_token (parser->lexer); + it.name.id = tok->u.value; + it.name.loc = tok->location; + cp_lexer_consume_token (parser->lexer); + + if (cp_lexer_next_token_is (parser->lexer, CPP_OPEN_PAREN)) + { + matching_parens parens2; + parens2.consume_open (parser); + + if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)) + { + tok = cp_lexer_peek_token (parser->lexer); + it.arg.id = tok->u.value; + it.arg.loc = tok->location; + cp_lexer_consume_token (parser->lexer); + } + else + { + cp_parser_error (parser, "expected identifier"); + cp_parser_skip_to_closing_parenthesis (parser, + /*recovering=*/true, + /*or_comma=*/false, + /*consume_paren=*/true); + goto end; + } + cp_parser_skip_to_closing_parenthesis (parser, + /*recovering=*/false, + /*or_comma=*/false, + /*consume_paren=*/true); + } + } + + cur_list->safe_push (it); + + if (cp_lexer_next_token_is (parser->lexer, CPP_COMMA)) + cp_lexer_consume_token (parser->lexer); + else if (cp_lexer_next_token_is (parser->lexer, CPP_COLON)) + { + if (modifiers) + { + cp_parser_error (parser, "expected %<)%>"); + cp_parser_skip_to_closing_parenthesis (parser, + /*recovering=*/true, + /*or_comma=*/false, + /*consume_paren=*/true); + goto end; + } + else + { + cp_lexer_consume_token (parser->lexer); + modifiers = cur_list; + cur_list = new auto_vec (4); + } + } + else if (cp_lexer_next_token_is (parser->lexer, CPP_CLOSE_PAREN)) + { + gcc_assert (allocators == NULL); + allocators = cur_list; + cur_list = NULL; + break; + } + else + { + cp_parser_error (parser, "expected %<)%>"); + cp_parser_skip_to_closing_parenthesis (parser, + /*recovering=*/true, + /*or_comma=*/false, + /*consume_paren=*/true); + goto end; + } + } + + if (modifiers) + for (unsigned i = 0; i < modifiers->length (); i++) + { + item& it = (*modifiers)[i]; + const char *p = IDENTIFIER_POINTER (it.name.id); + int strcmp_traits = 1, strcmp_memspace = 1; + + if ((strcmp_traits = strcmp ("traits", p)) == 0 + || (strcmp_memspace = strcmp ("memspace", p)) == 0) + { + if ((strcmp_traits == 0 && traits_var != NULL_TREE) + || (strcmp_memspace == 0 && memspace_expr != NULL_TREE)) + { + error_at (it.name.loc, "duplicate %qs modifier", p); + goto end; + } + t = cp_parser_lookup_name_simple (parser, it.arg.id, it.arg.loc); + if (t == error_mark_node) + { + cp_parser_name_lookup_error (parser, it.arg.id, t, NLE_NULL, + it.arg.loc); + } + else if (strcmp_memspace == 0) + memspace_expr = t; + else if (strcmp_traits == 0) + traits_var = t; + else + gcc_unreachable (); + } + else + { + error_at (it.name.loc, "unknown modifier %qE", it.name.id); + goto end; + } + } + + if (allocators) + { + if (modifiers) + { + if (allocators->length () > 1) + { + error_at ((*allocators)[1].name.loc, + "% clause only accepts a single " + "allocator when using modifiers"); + goto end; + } + else if ((*allocators)[0].arg.id) + { + error_at ((*allocators)[0].arg.loc, + "legacy %<%E(%E)%> traits syntax not allowed in " + "% clause when using modifiers", + (*allocators)[0].name.id, (*allocators)[0].arg.id); + goto end; + } + } + + for (unsigned i = 0; i < allocators->length (); i++) + { + item& it = (*allocators)[i]; + t = cp_parser_lookup_name_simple (parser, it.name.id, it.name.loc); + if (t == error_mark_node) + { + cp_parser_name_lookup_error (parser, it.name.id, t, NLE_NULL, + it.name.loc); + goto end; + } + else if (t != error_mark_node) + { + tree t2 = NULL_TREE; + if (it.arg.id) + { + t2 = cp_parser_lookup_name_simple (parser, it.arg.id, + it.arg.loc); + if (t2 == error_mark_node) + { + cp_parser_name_lookup_error (parser, it.arg.id, t2, + NLE_NULL, it.arg.loc); + goto end; + } + } + else + t2 = traits_var; + + tree c = build_omp_clause (clause_loc, + OMP_CLAUSE_USES_ALLOCATORS); + OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c) = t; + OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c) = memspace_expr; + OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c) = t2; + OMP_CLAUSE_CHAIN (c) = nl; + nl = c; + } + } + } + end: + if (cur_list) + delete cur_list; + if (modifiers) + delete modifiers; + if (allocators) + delete allocators; + cp_parser_skip_to_closing_parenthesis (parser, + /*recovering=*/false, + /*or_comma=*/false, + /*consume_paren=*/true); + return nl; +} + /* OpenMP 2.5: lastprivate ( variable-list ) @@ -41856,6 +42086,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; @@ -46427,7 +46661,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 8090c71809f..47796ab68c1 100644 --- a/gcc/cp/semantics.cc +++ b/gcc/cp/semantics.cc @@ -8029,6 +8029,101 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) break; } gcc_unreachable (); + case OMP_CLAUSE_USES_ALLOCATORS: + t = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c); + if (TREE_CODE (t) == FIELD_DECL) + { + sorry_at (OMP_CLAUSE_LOCATION (c), "class members not yet " + "supported in % clause"); + remove = true; + break; + } + t = convert_from_reference (t); + if (TREE_CODE (TREE_TYPE (t)) != ENUMERAL_TYPE + || strcmp (IDENTIFIER_POINTER (TYPE_IDENTIFIER (TREE_TYPE (t))), + "omp_allocator_handle_t") != 0) + { + error_at (OMP_CLAUSE_LOCATION (c), + "allocator must be of % type"); + remove = true; + break; + } + if (TREE_CODE (t) == CONST_DECL) + { + /* Currently for pre-defined allocators in libgomp, we do not + require additional init/fini inside target regions, so discard + such clauses. */ + remove = true; + + if (strcmp (IDENTIFIER_POINTER (DECL_NAME (t)), + "omp_null_allocator") == 0) + { + error_at (OMP_CLAUSE_LOCATION (c), + "% cannot be used in " + "% clause"); + break; + } + + if (OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c) + || OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c)) + { + error_at (OMP_CLAUSE_LOCATION (c), + "modifiers cannot be used with pre-defined " + "allocators"); + break; + } + } + t = OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c); + if (t != NULL_TREE + && (TREE_CODE (t) != CONST_DECL + || TREE_CODE (TREE_TYPE (t)) != ENUMERAL_TYPE + || strcmp (IDENTIFIER_POINTER (TYPE_IDENTIFIER (TREE_TYPE (t))), + "omp_memspace_handle_t") != 0)) + { + error_at (OMP_CLAUSE_LOCATION (c), "memspace modifier must be " + "constant enum of % type"); + remove = true; + break; + } + t = OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c); + if (t != NULL_TREE) + { + bool type_err = false; + + if (TREE_CODE (TREE_TYPE (t)) != ARRAY_TYPE + || DECL_SIZE (t) == NULL_TREE) + type_err = true; + else + { + tree elem_t = TREE_TYPE (TREE_TYPE (t)); + if (TREE_CODE (elem_t) != RECORD_TYPE + || strcmp (IDENTIFIER_POINTER (TYPE_IDENTIFIER (elem_t)), + "omp_alloctrait_t") != 0 + || !TYPE_READONLY (elem_t)) + type_err = true; + } + if (type_err) + { + error_at (OMP_CLAUSE_LOCATION (c), "traits array %qE must be of " + "% type", t); + remove = true; + } + else + { + tree cst_val = decl_constant_value (t); + if (cst_val == t) + { + error_at (OMP_CLAUSE_LOCATION (c), "traits array must be " + "of constant values"); + + remove = true; + } + } + } + if (remove) + break; + pc = &OMP_CLAUSE_CHAIN (c); + continue; case OMP_CLAUSE_DEPEND: case OMP_CLAUSE_AFFINITY: t = OMP_CLAUSE_DECL (c); diff --git a/gcc/fortran/trans-array.cc b/gcc/fortran/trans-array.cc index bbb81f40aa9..3329f8f3729 100644 --- a/gcc/fortran/trans-array.cc +++ b/gcc/fortran/trans-array.cc @@ -6501,10 +6501,7 @@ 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 82bbc41b388..9905c882d29 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -2758,7 +2758,12 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, gfc_init_se (&se, NULL); gfc_conv_expr (&se, n->u2.allocator); gfc_add_block_to_block (block, &se.pre); - allocator_ = gfc_evaluate_now (se.expr, block); + t = se.expr; + if (DECL_P (t) && se.post.head == NULL_TREE) + allocator_ = (POINTER_TYPE_P (TREE_TYPE (t)) + ? build_fold_indirect_ref (t): t); + else + allocator_ = gfc_evaluate_now (t, block); gfc_add_block_to_block (block, &se.post); } OMP_CLAUSE_ALLOCATE_ALLOCATOR (node) = allocator_; @@ -3938,13 +3943,36 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, } break; case OMP_LIST_USES_ALLOCATORS: - /* Ignore pre-defined allocators as no special treatment is needed. */ for (; n != NULL; n = n->next) - if (n->sym->attr.flavor == FL_VARIABLE) - break; - if (n != NULL) - sorry_at (input_location, "% clause with traits " - "and memory spaces"); + { + if (!n->sym->attr.referenced) + continue; + tree node = build_omp_clause (input_location, + OMP_CLAUSE_USES_ALLOCATORS); + tree t; + if (n->sym->attr.flavor == FL_VARIABLE) + t = gfc_get_symbol_decl (n->sym); + else + { + t = gfc_conv_mpz_to_tree (n->sym->value->value.integer, + n->sym->ts.kind); + t = fold_convert (ptr_type_node, t); + } + OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR(node) = t; + if (n->u.memspace_sym) + { + n->u.memspace_sym->attr.referenced = true; + OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (node) + = gfc_get_symbol_decl (n->u.memspace_sym); + } + if (n->u2.traits_sym) + { + n->u2.traits_sym->attr.referenced = true; + OMP_CLAUSE_USES_ALLOCATORS_TRAITS (node) + = gfc_get_symbol_decl (n->u2.traits_sym); + } + omp_clauses = gfc_trans_add_clause (node, omp_clauses); + } break; default: break; diff --git a/gcc/fortran/types.def b/gcc/fortran/types.def index 7a465c89c5f..a9374d68d78 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 d52d71b9b6b..08bd37e2e22 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -1381,18 +1381,46 @@ gimplify_bind_expr (tree *expr_p, gimple_seq *pre_p) dynamic_allocators clause is present in the same compilation unit. */ bool missing_dyn_alloc = false; - if (alloc == NULL_TREE - && ((omp_requires_mask & OMP_REQUIRES_DYNAMIC_ALLOCATORS) - == 0)) + if ((omp_requires_mask & OMP_REQUIRES_DYNAMIC_ALLOCATORS) == 0) { /* This comes too early for omp_discover_declare_target..., but should at least catch the most common cases. */ missing_dyn_alloc - = cgraph_node::get (current_function_decl)->offloadable; + = (alloc == NULL_TREE + && cgraph_node::get (current_function_decl)->offloadable); for (struct gimplify_omp_ctx *ctx2 = ctx; ctx2 && !missing_dyn_alloc; ctx2 = ctx2->outer_context) if (ctx2->code == OMP_TARGET) - missing_dyn_alloc = true; + { + if (alloc == NULL_TREE) + missing_dyn_alloc = true; + else if (TREE_CODE (alloc) != INTEGER_CST) + { + tree alloc2 = alloc; + if (TREE_CODE (alloc2) == MEM_REF + || TREE_CODE (alloc2) == INDIRECT_REF) + alloc2 = TREE_OPERAND (alloc2, 0); + tree c2; + for (c2 = ctx2->clauses; c2; + c2 = OMP_CLAUSE_CHAIN (c2)) + if (OMP_CLAUSE_CODE (c2) + == OMP_CLAUSE_USES_ALLOCATORS) + { + tree t2 + = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c2); + if (operand_equal_p (alloc2, t2)) + break; + } + if (c2 == NULL_TREE) + error_at (EXPR_LOC_OR_LOC ( + alloc, DECL_SOURCE_LOCATION (t)), + "%qE in % clause inside a " + "target region must be specified in an " + "% clause on the " + "% directive", alloc2); + } + break; + } } if (missing_dyn_alloc) error_at (DECL_SOURCE_LOCATION (t), @@ -12240,6 +12268,21 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, nowait = 1; break; + case OMP_CLAUSE_USES_ALLOCATORS: + if (TREE_CODE (OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c)) + != INTEGER_CST) + { + decl = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c); + omp_add_variable (ctx, decl, GOVD_SEEN | GOVD_PRIVATE); + + decl = OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c); + if (decl && !DECL_INITIAL (decl)) + omp_add_variable (ctx, decl, GOVD_SEEN | GOVD_FIRSTPRIVATE); + } + else + remove = true; + break; + case OMP_CLAUSE_ORDERED: case OMP_CLAUSE_UNTIED: case OMP_CLAUSE_COLLAPSE: @@ -12375,6 +12418,49 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, remove = true; break; } + if ((omp_requires_mask & OMP_REQUIRES_DYNAMIC_ALLOCATORS) == 0 + && OMP_CLAUSE_ALLOCATE_ALLOCATOR (c) + && TREE_CODE (OMP_CLAUSE_ALLOCATE_ALLOCATOR (c)) != INTEGER_CST) + { + tree allocator = OMP_CLAUSE_ALLOCATE_ALLOCATOR (c); + tree clauses = NULL_TREE; + + /* Get clause list of the nearest enclosing target construct. */ + if (ctx->code == OMP_TARGET) + clauses = *orig_list_p; + else + { + struct gimplify_omp_ctx *tctx = ctx->outer_context; + while (tctx && tctx->code != OMP_TARGET) + tctx = tctx->outer_context; + if (tctx) + clauses = tctx->clauses; + } + + if (clauses) + { + tree uc; + if (TREE_CODE (allocator) == MEM_REF + || TREE_CODE (allocator) == INDIRECT_REF) + allocator = TREE_OPERAND (allocator, 0); + for (uc = clauses; uc; uc = OMP_CLAUSE_CHAIN (uc)) + if (OMP_CLAUSE_CODE (uc) == OMP_CLAUSE_USES_ALLOCATORS) + { + tree uc_allocator + = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (uc); + if (operand_equal_p (allocator, uc_allocator)) + break; + } + if (uc == NULL_TREE) + { + error_at (OMP_CLAUSE_LOCATION (c), "allocator %qE " + "requires % clause in " + "target region", allocator, allocator); + remove = true; + break; + } + } + } if (gimplify_expr (&OMP_CLAUSE_ALLOCATE_ALLOCATOR (c), pre_p, NULL, is_gimple_val, fb_rvalue) == GS_ERROR) { @@ -13411,6 +13497,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, case OMP_CLAUSE_FINALIZE: case OMP_CLAUSE_INCLUSIVE: case OMP_CLAUSE_EXCLUSIVE: + case OMP_CLAUSE_USES_ALLOCATORS: break; case OMP_CLAUSE_NOHOST: @@ -15777,6 +15864,92 @@ 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 = NULL_TREE; + tree omp_destroy_allocator_fn = NULL_TREE; + + for (tree *cp = &OMP_CLAUSES (expr); *cp != NULL; + cp = &OMP_CLAUSE_CHAIN (*cp)) + 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); + + if (omp_init_allocator_fn == NULL_TREE) + { + omp_init_allocator_fn + = builtin_decl_explicit (BUILT_IN_OMP_INIT_ALLOCATOR); + omp_destroy_allocator_fn + = builtin_decl_explicit (BUILT_IN_OMP_DESTROY_ALLOCATOR); + } + tree ntraits, traits_var; + if (traits == NULL_TREE) + { + ntraits = integer_zero_node; + traits_var = null_pointer_node; + } + else if (DECL_INITIAL (traits)) + { + location_t loc = OMP_CLAUSE_LOCATION (c); + tree t = DECL_INITIAL (traits); + gcc_assert (TREE_CODE (t) == CONSTRUCTOR); + ntraits = build_int_cst (integer_type_node, + CONSTRUCTOR_NELTS (t)); + t = get_initialized_tmp_var (t, &init_seq, NULL); + traits_var = build_fold_addr_expr_loc (loc, t); + } + else + { + location_t loc = OMP_CLAUSE_LOCATION (c); + gcc_assert (TREE_CODE (TREE_TYPE (traits)) == ARRAY_TYPE); + tree t = TYPE_DOMAIN (TREE_TYPE (traits)); + tree min = TYPE_MIN_VALUE (t); + tree max = TYPE_MAX_VALUE (t); + gcc_assert (TREE_CODE (min) == INTEGER_CST + && TREE_CODE (max) == INTEGER_CST); + t = fold_build2_loc (loc, MINUS_EXPR, TREE_TYPE (min), + max, min); + t = fold_build2_loc (loc, PLUS_EXPR, TREE_TYPE (min), + t, build_int_cst (TREE_TYPE (min), 1)); + ntraits = t; + traits_var = build_fold_addr_expr_loc (loc, traits); + } + + if (memspace == NULL_TREE) + memspace = build_int_cst (pointer_sized_int_node, 0); + else + memspace = fold_convert (pointer_sized_int_node, + memspace); + + tree call = build_call_expr_loc (OMP_CLAUSE_LOCATION (c), + omp_init_allocator_fn, 3, + memspace, ntraits, + traits_var); + call = fold_convert (TREE_TYPE (allocator), call); + gimplify_assign (allocator, call, &init_seq); + + call = build_call_expr_loc (OMP_CLAUSE_LOCATION (c), + omp_destroy_allocator_fn, 1, + allocator); + gimplify_and_add (call, &fini_seq); + } + + 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); + body = bind; + } + } } else gimplify_and_add (OMP_BODY (expr), &body); diff --git a/gcc/omp-builtins.def b/gcc/omp-builtins.def index ed78d49d205..c07f8533c0a 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/omp-low.cc b/gcc/omp-low.cc index 161bcfeec05..a53a620fc1f 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -1166,6 +1166,36 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) && omp_maybe_offloaded_ctx (ctx)) error_at (OMP_CLAUSE_LOCATION (c), "% clause must" " specify an allocator here"); + if ((omp_requires_mask & OMP_REQUIRES_DYNAMIC_ALLOCATORS) == 0 + && OMP_CLAUSE_ALLOCATE_ALLOCATOR (c) != NULL_TREE + && DECL_P (OMP_CLAUSE_ALLOCATE_ALLOCATOR (c)) + && !DECL_ARTIFICIAL (OMP_CLAUSE_ALLOCATE_ALLOCATOR (c))) + { + tree alloc2 = OMP_CLAUSE_ALLOCATE_ALLOCATOR (c); + if (TREE_CODE (alloc2) == MEM_REF + || TREE_CODE (alloc2) == INDIRECT_REF) + alloc2 = TREE_OPERAND (alloc2, 0); + omp_context *ctx2 = ctx; + for (; ctx2; ctx2 = ctx2->outer) + if (is_gimple_omp_offloaded (ctx2->stmt)) + break; + if (ctx2 != NULL) + { + tree c2 = gimple_omp_target_clauses (ctx2->stmt); + for (; c2; c2 = OMP_CLAUSE_CHAIN (c2)) + if (OMP_CLAUSE_CODE (c2) == OMP_CLAUSE_USES_ALLOCATORS + && operand_equal_p ( + alloc2, OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c2))) + break; + if (c2 == NULL_TREE) + error_at (EXPR_LOC_OR_LOC (OMP_CLAUSE_ALLOCATE_ALLOCATOR (c), + OMP_CLAUSE_LOCATION (c)), + "allocator %qE in % clause inside a " + "target region must be specified in an " + "% clause on the % " + "directive", alloc2); + } + } if (ctx->allocate_map == NULL) ctx->allocate_map = new hash_map; tree val = integer_zero_node; @@ -1756,6 +1786,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_FINALIZE: case OMP_CLAUSE_TASK_REDUCTION: case OMP_CLAUSE_ALLOCATE: + case OMP_CLAUSE_USES_ALLOCATORS: break; case OMP_CLAUSE_ALIGNED: @@ -1974,6 +2005,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_FINALIZE: case OMP_CLAUSE_FILTER: case OMP_CLAUSE__CONDTEMP_: + case OMP_CLAUSE_USES_ALLOCATORS: break; case OMP_CLAUSE__CACHE_: 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..5a2e4a90e54 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/uses_allocators-1.c @@ -0,0 +1,46 @@ +typedef enum omp_allocator_handle_t +#if __cplusplus >= 201103L +: __UINTPTR_TYPE__ +#endif +{ + omp_default_mem_alloc = 1, + omp_low_lat_mem_alloc = 5, + __omp_allocator_handle_t_max__ = __UINTPTR_MAX__ +} omp_allocator_handle_t; + +typedef struct omp_alloctrait_t +{ + int key; + int value; +} omp_alloctrait_t; + +extern void *omp_alloc (__SIZE_TYPE__, omp_allocator_handle_t); + +void +f (omp_allocator_handle_t my_alloc) +{ + #pragma omp target + { + int a; /* { dg-error "'my_alloc' in 'allocator' clause inside a target region must be specified in an 'uses_allocators' clause on the 'target' directive" "" { target c } } */ + #pragma omp allocate(a) allocator(my_alloc) /* { dg-message "sorry, unimplemented: '#pragma omp allocate' not yet supported" "" { target c++ } } */ + a = 5; + void *prt = omp_alloc(32, my_alloc); + #pragma omp parallel allocate(allocator(my_alloc) : a) firstprivate(a) /* { dg-error "allocator 'my_alloc' in 'allocate' clause inside a target region must be specified in an 'uses_allocators' clause on the 'target' directive" } */ + a = 7; + } +} + +void +g (omp_allocator_handle_t my_alloc) +{ + /* The following defines a default-mem-space allocator with no extra traits. */ + #pragma omp target uses_allocators(my_alloc) + { + int a; + #pragma omp allocate(a) allocator(my_alloc) /* { dg-message "sorry, unimplemented: '#pragma omp allocate' not yet supported" "" { target c++ } } */ + a = 5; + void *prt = omp_alloc(32, my_alloc); + #pragma omp parallel allocate(allocator(my_alloc) : a) firstprivate(a) + a = 7; + } +} 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..4dd1f13100a --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/uses_allocators-2.c @@ -0,0 +1,33 @@ +typedef enum omp_allocator_handle_t +#if __cplusplus >= 201103L +: __UINTPTR_TYPE__ +#endif +{ + omp_default_mem_alloc = 1, + omp_low_lat_mem_alloc = 5, + __omp_allocator_handle_t_max__ = __UINTPTR_MAX__ +} omp_allocator_handle_t; + +typedef struct omp_alloctrait_t +{ + int key; + int value; +} omp_alloctrait_t; + +void +f () +{ + omp_alloctrait_t trait[1] = {{1,1}}; + omp_allocator_handle_t my_alloc; + #pragma omp target uses_allocators(traits(trait) : my_alloc) /* { dg-error "traits array 'trait' must be of 'const omp_alloctrait_t \\\[\\\]' type" } */ + ; +} + +void +g () +{ + const omp_alloctrait_t trait[1] = {{1,1}}; + omp_allocator_handle_t my_alloc; + #pragma omp target uses_allocators(traits(trait) : my_alloc) + ; +} diff --git a/gcc/testsuite/gfortran.dg/gomp/allocate-1.f90 b/gcc/testsuite/gfortran.dg/gomp/allocate-1.f90 index 8bc6b768778..0463f0e0af9 100644 --- a/gcc/testsuite/gfortran.dg/gomp/allocate-1.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/allocate-1.f90 @@ -24,6 +24,10 @@ module omp_lib_kinds parameter :: omp_pteam_mem_alloc = 7 integer (kind=omp_allocator_handle_kind), & parameter :: omp_thread_mem_alloc = 8 + + integer, parameter :: omp_memspace_handle_kind = c_intptr_t + integer (omp_memspace_handle_kind), & + parameter :: omp_default_mem_space = 0 end module subroutine bar (a, b, c) @@ -80,7 +84,8 @@ subroutine foo(x, y) !$omp target teams distribute parallel do private (x) firstprivate (y) & !$omp allocate ((omp_default_mem_alloc + 0):z) allocate & - !$omp (omp_default_mem_alloc: x, y) allocate (h: r) lastprivate (z) reduction(+:r) + !$omp (omp_default_mem_alloc: x, y) allocate (h: r) lastprivate (z) reduction(+:r) & + !$omp uses_allocators(memspace(omp_default_mem_space) : h) do i = 1, 10 call bar (0, x, z); call bar2 (1, y, r); diff --git a/gcc/testsuite/gfortran.dg/gomp/scope-6.f90 b/gcc/testsuite/gfortran.dg/gomp/scope-6.f90 index 4c4f5e034f7..39a65904c33 100644 --- a/gcc/testsuite/gfortran.dg/gomp/scope-6.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/scope-6.f90 @@ -20,4 +20,4 @@ contains end end -! { dg-final { scan-tree-dump "omp scope private\\(a\\) firstprivate\\(b\\) reduction\\(\\+:c\\) allocate\\(allocator\\(D\\.\[0-9\]+\\):a\\) allocate\\(allocator\\(D\\.\[0-9\]+\\):b\\) allocate\\(allocator\\(D\\.\[0-9\]+\\):c\\)" "original" } } +! { dg-final { scan-tree-dump "omp scope private\\(a\\) firstprivate\\(b\\) reduction\\(\\+:c\\) allocate\\(allocator\\(h\\):a\\) allocate\\(allocator\\(h\\):b\\) allocate\\(allocator\\(h\\):c\\)" "original" } } 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..b20515b3691 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-1.f90 @@ -0,0 +1,23 @@ +use iso_c_binding +implicit none + integer, parameter :: omp_allocator_handle_kind = c_intptr_t + integer, parameter :: omp_alloctrait_key_kind = c_int + integer, parameter :: omp_alloctrait_val_kind = c_intptr_t + type omp_alloctrait + integer (kind=omp_alloctrait_key_kind) key + integer (kind=omp_alloctrait_val_kind) value + end type omp_alloctrait +contains +subroutine x +integer :: mem +type(omp_alloctrait), parameter:: mem2(1) = [omp_alloctrait(1,1)] +integer(omp_allocator_handle_kind) :: var +!$omp target uses_allocators(memspace(omp_default_mem_space), traits(mem2) : var) defaultmap(none) +block; +type(c_ptr) ::c +c = omp_alloc(omp_default_mem_space, 20_8) +end block +!$omp target uses_allocators(omp_default_mem_alloc, var(mem2)) +block; end block +end +end diff --git a/gcc/tree-core.h b/gcc/tree-core.h index 65e51b939a2..c76e7d6116f 100644 --- a/gcc/tree-core.h +++ b/gcc/tree-core.h @@ -533,6 +533,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 1fadd752d05..8cf479d8508 100644 --- a/gcc/tree-pretty-print.cc +++ b/gcc/tree-pretty-print.cc @@ -800,6 +800,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 a3d907a4ac0..46f41cb7167 100644 --- a/gcc/tree.cc +++ b/gcc/tree.cc @@ -328,6 +328,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[] = @@ -421,6 +422,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 086b55f0375..88befefe110 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -1958,6 +1958,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) diff --git a/libgomp/testsuite/libgomp.c++/c++.exp b/libgomp/testsuite/libgomp.c++/c++.exp index ed096e17b9c..5be949bb611 100644 --- a/libgomp/testsuite/libgomp.c++/c++.exp +++ b/libgomp/testsuite/libgomp.c++/c++.exp @@ -1,6 +1,15 @@ load_lib libgomp-dg.exp load_gcc_lib gcc-dg.exp +proc check_effective_target_c { } { + return 0 +} + +proc check_effective_target_c++ { } { + return 1 +} + + if { $blddir != "" } { set libstdc++_library_path "../libstdc++-v3/src/.libs" set shlib_ext [get_shlib_extension] diff --git a/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-1.c b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-1.c new file mode 100644 index 00000000000..eaca686d037 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-1.c @@ -0,0 +1,53 @@ +/* { 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 uses_allocators\\(bar: memspace\\(\\), traits\\(\\)\\)" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(foo: memspace\\(\\), traits\\(foo_traits\\)\\)" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(bar: memspace\\(\\), traits\\(foo_traits\\)\\) uses_allocators\\(foo: memspace\\(\\), traits\\(foo_traits\\)\\)" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(foo: memspace\\(omp_high_bw_mem_space\\), traits\\(\\)\\)" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(bar: memspace\\(\\), traits\\(foo_traits\\)\\)" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(bar: memspace\\(omp_high_bw_mem_space\\), traits\\(foo_traits\\)\\)" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(bar: memspace\\(omp_high_bw_mem_space\\), traits\\(foo_traits\\)\\) uses_allocators\\(foo: memspace\\(\\), traits\\(\\)\\)" "original" } } */ + +/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(bar: memspace\\(\\), traits\\(\\)\\) private\\(bar\\)" "gimple" } } */ +/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(foo: memspace\\(\\), traits\\(foo_traits\\)\\) private\\(foo\\)" "gimple" } } */ +/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(bar: memspace\\(\\), traits\\(foo_traits\\)\\) uses_allocators\\(foo: memspace\\(\\), traits\\(foo_traits\\)\\) private\\(bar\\) private\\(foo\\)" "gimple" } } */ +/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(foo: memspace\\(omp_high_bw_mem_space\\), traits\\(\\)\\) private\\(foo\\)" "gimple" } } */ +/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(bar: memspace\\(\\), traits\\(foo_traits\\)\\) private\\(bar\\)" "gimple" } } */ +/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(bar: memspace\\(omp_high_bw_mem_space\\), traits\\(foo_traits\\)\\) private\\(bar\\)" "gimple" } } */ +/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(bar: memspace\\(omp_high_bw_mem_space\\), traits\\(foo_traits\\)\\) uses_allocators\\(foo: memspace\\(\\), traits\\(\\)\\) private\\(bar\\) private\\(foo\\)" "gimple" } } */ + +/* { 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/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-2.c b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-2.c new file mode 100644 index 00000000000..f350c0a409e --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-2.c @@ -0,0 +1,39 @@ +/* { dg-do compile } */ + +#include + +omp_alloctrait_key_t k; +omp_alloctrait_value_t v; + +int main (void) +{ + omp_allocator_handle_t foo, bar; + const omp_alloctrait_t traits_array[] = { { omp_atk_pinned, omp_atv_true }, + { omp_atk_partition, omp_atv_nearest } }; + + #pragma omp target uses_allocators (baz) /* { dg-error "'baz' undeclared .first use in this function." "" { target c } } */ + ; /* { dg-error "'baz' has not been declared" "" { target c++ } .-1 } */ + #pragma omp target uses_allocators (foo (xyz)) /* { dg-error "'xyz' undeclared .first use in this function." "" { target c } } */ + ; /* { dg-error "'xyz' has not been declared" "" { target c++ } .-1 } */ + #pragma omp target uses_allocators (foo (traits_array), baz (traits_array)) /* { dg-error "'baz' has not been declared" "" { target c++ } } */ + ; + #pragma omp target uses_allocators (memspace(omp_no_such_space) : foo) /* { dg-error "'omp_no_such_space' undeclared .first use in this function." "" { target c } } */ + ; /* { dg-error "'omp_no_such_space' has not been declared" "" { target c++ } .-1 } */ + #pragma omp target uses_allocators (memspace(1) : foo) /* { dg-error "expected identifier before numeric constant" } */ + ; /* { dg-error "expected '\\\)' before ':' token" "" { target c } .-1 } */ + #pragma omp target uses_allocators (memspace(omp_no_such_space) : foo, bar) /* { dg-error "'uses_allocators' clause only accepts a single allocator when using modifiers" } */ + ; /* { dg-error "'omp_no_such_space' has not been declared" "" { target c++ } .-1 } */ + #pragma omp target uses_allocators (traits(xyz) : bar) /* { dg-error "traits array must be of 'const omp_alloctrait_t \\\[\\\]' type" "" { target c } } */ + ; /* { dg-error "'xyz' has not been declared" "" { target c++ } .-1 } */ + #pragma omp target uses_allocators (memspace(omp_high_bw_mem_space), traits(traits_array), memspace (omp_no_such_space) : bar) /* { dg-error "duplicate 'memspace' modifier" } */ + ; + #pragma omp target uses_allocators (traitz(traits_array), memspace(omp_high_bw_mem_space) : bar) /* { dg-error "unknown modifier 'traitz'" } */ + ; + #pragma omp target uses_allocators (omp_null_allocator) /* { dg-error "'omp_null_allocator' cannot be used in 'uses_allocators' clause" } */ + ; + #pragma omp target uses_allocators (memspace(omp_high_bw_mem_space) : foo, bar) /* { dg-error "'uses_allocators' clause only accepts a single allocator when using modifiers" } */ + ; + #pragma omp target uses_allocators (memspace(omp_high_bw_mem_space) : foo(foo_traits)) /* { dg-error "legacy 'foo\\\(foo_traits\\\)' traits syntax not allowed in 'uses_allocators' clause when using modifiers" } */ + ; + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-3.c b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-3.c new file mode 100644 index 00000000000..de9ab9213aa --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-3.c @@ -0,0 +1,37 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" } */ + +#include + +int main (void) +{ + omp_allocator_handle_t memspace, traits; + const omp_alloctrait_t mytraits[] = { { omp_atk_pinned, omp_atv_true }, + { omp_atk_partition, omp_atv_nearest } }; + #pragma omp target uses_allocators (memspace) + ; + #pragma omp target uses_allocators (traits) + ; + #pragma omp target uses_allocators (traits, memspace) + ; + #pragma omp target uses_allocators (traits (mytraits)) + ; + #pragma omp target uses_allocators (memspace (mytraits), omp_default_mem_alloc) + ; + return 0; +} + +/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(memspace: memspace\\(\\), traits\\(\\)\\)" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(traits: memspace\\(\\), traits\\(\\)\\)" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(memspace: memspace\\(\\), traits\\(\\)\\) uses_allocators\\(traits: memspace\\(\\), traits\\(\\)\\)" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(traits: memspace\\(\\), traits\\(mytraits\\)\\)" "original" } } */ +/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(memspace: memspace\\(\\), traits\\(mytraits\\)\\)" "original" } } */ + +/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(memspace: memspace\\(\\), traits\\(\\)\\) private\\(memspace\\)" "gimple" } } */ +/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(traits: memspace\\(\\), traits\\(\\)\\) private\\(traits\\)" "gimple" } } */ +/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(memspace: memspace\\(\\), traits\\(\\)\\) uses_allocators\\(traits: memspace\\(\\), traits\\(\\)\\) private\\(traits\\) private\\(memspace\\)" "gimple" } } */ +/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(traits: memspace\\(\\), traits\\(mytraits\\)\\) private\\(traits\\)" "gimple" } } */ +/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(memspace: memspace\\(\\), traits\\(mytraits\\)\\) private\\(memspace\\)" "gimple" } } */ + +/* { 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/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-4.c b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-4.c new file mode 100644 index 00000000000..fabf2a8a319 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-4.c @@ -0,0 +1,53 @@ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +#include +#include + +int +main () +{ + int x, *xbuf[10]; + omp_allocator_handle_t my_alloc; + const omp_alloctrait_t trait[1]= {{omp_atk_alignment,128}}; + + #pragma omp target uses_allocators(omp_low_lat_mem_alloc) map(tofrom: x, xbuf) defaultmap(none) + #pragma omp parallel allocate(allocator(omp_low_lat_mem_alloc), align(128): x, xbuf) if(0) firstprivate(x, xbuf) + { + if ((uintptr_t) &x % 128 != 0) + __builtin_abort (); + if ((uintptr_t) xbuf % 128 != 0) + __builtin_abort (); + } + + my_alloc = (omp_allocator_handle_t) 0xABCD; + + #pragma omp target uses_allocators(traits(trait): my_alloc) defaultmap(none) map(tofrom: x, xbuf) + #pragma omp parallel allocate(allocator(my_alloc): x, xbuf) if(0) firstprivate(x, xbuf) + { + if ((uintptr_t) &x % 128 != 0) + __builtin_abort (); + if ((uintptr_t) xbuf % 128 != 0) + __builtin_abort (); + } + + if (my_alloc != (omp_allocator_handle_t) 0xABCD) + __builtin_abort (); + + /* The following creates an allocator with empty traits + default mem space. */ + #pragma omp target uses_allocators(my_alloc) map(tofrom: x, xbuf) defaultmap(none) + #pragma omp parallel allocate(allocator(my_alloc), align(128): x, xbuf) if(0) firstprivate(x, xbuf) + { + if ((uintptr_t) &x % 128 != 0) + __builtin_abort (); + if ((uintptr_t) xbuf % 128 != 0) + __builtin_abort (); + } + + if (my_alloc != (omp_allocator_handle_t) 0xABCD) + __builtin_abort (); + + return 0; +} + +/* { dg-final { scan-tree-dump-times "#pragma omp target .*private\\(my_alloc\\).*uses_allocators\\(my_alloc: memspace\\(\\), traits\\(trait\\)\\)" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "#pragma omp target .*private\\(my_alloc\\).*uses_allocators\\(my_alloc: memspace\\(\\), traits\\(\\)\\)" 1 "gimple" } } */ diff --git a/libgomp/testsuite/libgomp.c/c.exp b/libgomp/testsuite/libgomp.c/c.exp index aae282478db..4b59957d1f3 100644 --- a/libgomp/testsuite/libgomp.c/c.exp +++ b/libgomp/testsuite/libgomp.c/c.exp @@ -3,6 +3,14 @@ load_gcc_lib gcc-dg.exp lappend ALWAYS_CFLAGS "compiler=$GCC_UNDER_TEST" +proc check_effective_target_c { } { + return 1 +} + +proc check_effective_target_c++ { } { + return 0 +} + # If a testcase doesn't have special options, use these. if ![info exists DEFAULT_CFLAGS] then { set DEFAULT_CFLAGS "-O2" diff --git a/libgomp/testsuite/libgomp.fortran/uses_allocators_2.f90 b/libgomp/testsuite/libgomp.fortran/uses_allocators_2.f90 index 07327969775..bb984033413 100644 --- a/libgomp/testsuite/libgomp.fortran/uses_allocators_2.f90 +++ b/libgomp/testsuite/libgomp.fortran/uses_allocators_2.f90 @@ -3,8 +3,6 @@ ! Minimal test for valid code: ! - predefined allocators do not need any special treatment in uses_allocators ! (as 'requires dynamic_allocators' is the default). -! -! - Non-predefined allocators are currently rejected ('sorry)' subroutine test use omp_lib @@ -35,22 +33,22 @@ subroutine non_predef integer(kind=omp_allocator_handle_kind) :: a1, a2, a3 - !$omp target uses_allocators(omp_default_mem_alloc, a1(trait), a2(trait2)) ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" } + !$omp target uses_allocators(omp_default_mem_alloc, a1(trait), a2(trait2)) block; end block - !$omp target parallel uses_allocators(omp_default_mem_alloc, a1(trait), a2(trait2)) ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" } + !$omp target parallel uses_allocators(omp_default_mem_alloc, a1(trait), a2(trait2)) block; end block !$omp target uses_allocators(traits(trait):a1) & - !$omp& uses_allocators ( memspace ( omp_low_lat_mem_space ) , traits ( trait2 ) : a2 , a3) ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" } + !$omp& uses_allocators ( memspace ( omp_low_lat_mem_space ) , traits ( trait2 ) : a2 , a3) block; end block !$omp target parallel uses_allocators(traits(trait):a1) & - !$omp& uses_allocators ( memspace ( omp_low_lat_mem_space ) , traits ( trait2 ) : a2 , a3) ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" } + !$omp& uses_allocators ( memspace ( omp_low_lat_mem_space ) , traits ( trait2 ) : a2 , a3) block; end block - !$omp target uses_allocators ( traits(trait2) , memspace ( omp_low_lat_mem_space ) : a2 , a3) ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" } + !$omp target uses_allocators ( traits(trait2) , memspace ( omp_low_lat_mem_space ) : a2 , a3) block; end block end subroutine @@ -62,7 +60,7 @@ subroutine trait_present integer(kind=omp_allocator_handle_kind) :: a1 ! Invalid in OpenMP 5.0 / 5.1, but valid since 5.2 the same as omp_default_mem_space + emptry traits array - !$omp target uses_allocators ( a1 ) ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" } + !$omp target uses_allocators ( a1 ) block; end block end @@ -76,13 +74,13 @@ subroutine odd_names integer(kind=omp_allocator_handle_kind) :: traits integer(kind=omp_allocator_handle_kind) :: memspace - !$omp target uses_allocators ( traits(trait1), memspace(trait1) ) ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" } + !$omp target uses_allocators ( traits(trait1), memspace(trait1) ) block; end block - !$omp target uses_allocators ( traits(trait1), memspace(omp_low_lat_mem_space) : traits) ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" } + !$omp target uses_allocators ( traits(trait1), memspace(omp_low_lat_mem_space) : traits) block; end block - !$omp target uses_allocators ( memspace(omp_low_lat_mem_space), traits(trait1) : memspace) ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" } + !$omp target uses_allocators ( memspace(omp_low_lat_mem_space), traits(trait1) : memspace) block; end block end @@ -94,6 +92,6 @@ subroutine more_checks integer(kind=omp_allocator_handle_kind) :: a1, a2(4) integer(kind=1) :: a3 - !$omp target uses_allocators(memspace (omp_low_lat_mem_space) : a1 ) ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" } + !$omp target uses_allocators(memspace (omp_low_lat_mem_space) : a1 ) block; end block end diff --git a/libgomp/testsuite/libgomp.fortran/uses_allocators_3.f90 b/libgomp/testsuite/libgomp.fortran/uses_allocators_3.f90 new file mode 100644 index 00000000000..8acdd4223f9 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/uses_allocators_3.f90 @@ -0,0 +1,62 @@ +! { 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 + if (foo == 0) stop 1 + 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\\) uses_allocators\\(bar: memspace\\(\\), traits\\(\\)\\)" "original" } } +! { dg-final { scan-tree-dump "pragma omp target" "original" } } +! { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(bar: memspace\\(\\), traits\\(traits_array\\)\\) uses_allocators\\(foo: memspace\\(\\), traits\\(traits_array\\)\\)" "original" } } +! { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(bar: memspace\\(\\), traits\\(traits_array\\)\\)" "original" } } +! { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(bar: memspace\\(omp_low_lat_mem_space\\), traits\\(\\)\\)" "original" } } +! { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(bar: memspace\\(omp_high_bw_mem_space\\), traits\\(traits_array\\)\\)" "original" } } + +! { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) allocate\\(allocator\\(bar\\):arr\\) uses_allocators\\(bar: memspace\\(\\), traits\\(\\)\\) private\\(bar\\)" "gimple" } } +! { dg-final { scan-tree-dump "pragma omp target" "gimple" } } +! { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(bar: memspace\\(\\), traits\\(traits_array\\)\\) uses_allocators\\(foo: memspace\\(\\), traits\\(traits_array\\)\\) private\\(foo\\) private\\(bar\\)" "gimple" } } +! { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(bar: memspace\\(\\), traits\\(traits_array\\)\\) private\\(bar\\)" "gimple" } } +! { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(bar: memspace\\(omp_low_lat_mem_space\\), traits\\(\\)\\) firstprivate\\(omp_low_lat_mem_space\\) private\\(bar\\)" "gimple" } } +! { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(bar: memspace\\(omp_high_bw_mem_space\\), traits\\(traits_array\\)\\) firstprivate\\(omp_high_bw_mem_space\\) private\\(bar\\)" "gimple" } } + +! { 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/libgomp/testsuite/libgomp.fortran/uses_allocators_4.f90 b/libgomp/testsuite/libgomp.fortran/uses_allocators_4.f90 new file mode 100644 index 00000000000..00f1dcb2763 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/uses_allocators_4.f90 @@ -0,0 +1,54 @@ +! { 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 "Allocator 'omp_non_existant_alloc' at .1. in USES_ALLOCATORS must be a scalar integer of kind 'omp_allocator_handle_kind'" } + block ! { dg-error "Symbol 'omp_non_existant_alloc' at .1. has no IMPLICIT type; did you mean 'omp_const_mem_alloc'\?" "" { target *-*-* } .-1 } + end block + + !$omp target uses_allocators(bar(traits_array), foo (traits_array), ) ! { dg-error "Invalid character in name" } + block + end block + + !$omp target uses_allocators(traits(xyz) : bar) ! { dg-error "Symbol 'xyz' at .1. has no IMPLICIT type" } + block ! { dg-error "Traits array 'xyz' in USES_ALLOCATORS .1. must be a one-dimensional named constant array of type 'omp_alloctrait'" "" { target *-*-* } .-1 } + end block + + !$omp target uses_allocators(memspace(omp_non_existant_mem_space) : foo) ! { dg-error "Symbol 'omp_non_existant_mem_space' at .1. has no IMPLICIT type; did you mean 'omp_const_mem_space'\?" } + ! { dg-error "Memspace 'omp_non_existant_mem_space' at .1. in USES_ALLOCATORS must be a predefined memory space" "" { target *-*-* } .-1 } + + block + end block + + !$omp target uses_allocators(traits(traits_array), traits(traits_array) : bar) ! { dg-error "Duplicate TRAITS modifier at .1. in USES_ALLOCATORS clause" } + block + end block + + !$omp target uses_allocators(memspace(omp_default_mem_space), memspace(omp_default_mem_space) : foo) ! { dg-error "Duplicate MEMSPACE modifier at .1. in USES_ALLOCATORS clause" } + block + end block + + !$omp target uses_allocators(memspace(omp_default_mem_space), traits(traits_array), traits(traits_array) : foo) ! { dg-error "Duplicate TRAITS modifier at .1. in USES_ALLOCATORS clause" } + block + end block + + !$omp target uses_allocators (omp_null_allocator) ! { dg-error "Allocator 'omp_null_allocator' at .1. in USES_ALLOCATORS must either a variable or a predefined allocator" } + block + end block + + !$omp target uses_allocators (memspace(omp_high_bw_mem_space) : foo, bar) + block + end block + + !$omp target uses_allocators (memspace(omp_high_bw_mem_space) : foo(foo_traits)) ! { dg-error "70:Unexpected '\\(' at .1." } + block + end block + +end program main diff --git a/libgomp/testsuite/libgomp.fortran/uses_allocators_5.f90 b/libgomp/testsuite/libgomp.fortran/uses_allocators_5.f90 new file mode 100644 index 00000000000..00f87109d2c --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/uses_allocators_5.f90 @@ -0,0 +1,14 @@ +! { dg-do compile } + +program main + use omp_lib + implicit none + integer, allocatable :: arr(:) + integer (omp_allocator_handle_kind) :: bar + + !$omp target allocate(bar : arr) ! { dg-error "allocator 'bar' requires 'uses_allocators.bar.' clause in target region" } + block + allocate(arr(100)) + end block + +end program main diff --git a/libgomp/testsuite/libgomp.fortran/uses_allocators_6.f90 b/libgomp/testsuite/libgomp.fortran/uses_allocators_6.f90 new file mode 100644 index 00000000000..3799e3cec73 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/uses_allocators_6.f90 @@ -0,0 +1,50 @@ +! { dg-additional-options "-fdump-tree-gimple" } + +program main + use iso_c_binding + use omp_lib + implicit none (type, external) + integer :: x, xbuf(10) + integer(c_intptr_t) :: iptr + integer(omp_allocator_handle_kind) :: my_alloc + type(omp_alloctrait), parameter :: trait(*) = [omp_alloctrait(omp_atk_alignment, 128)] + + !$omp target uses_allocators(omp_low_lat_mem_alloc) map(tofrom: x, xbuf) defaultmap(none) + !$omp parallel allocate(allocator(omp_low_lat_mem_alloc), align(128): x, xbuf) if(.false.) firstprivate(x, xbuf) + if (mod (TRANSFER (loc(x), iptr), 128) /= 0) & + stop 1 + if (mod (TRANSFER (loc(xbuf), iptr), 128) /= 0) & + stop 2 + !$omp end parallel + !$omp end target + + my_alloc = transfer(int(z'ABCD', omp_allocator_handle_kind), my_alloc) + + !$omp target uses_allocators(traits(trait): my_alloc) defaultmap(none) map(tofrom: x, xbuf) + !$omp parallel allocate(allocator(my_alloc): x, xbuf) if(.false.) firstprivate(x, xbuf) + if (mod (TRANSFER (loc(x), iptr), 128) /= 0) & + stop 3 + if (mod (TRANSFER (loc(xbuf), iptr), 128) /= 0) & + stop 4 + !$omp end parallel + !$omp end target + + if (transfer(my_alloc, 0_omp_allocator_handle_kind) /= int(z'ABCD', omp_allocator_handle_kind)) & + stop 5 + + ! The following creates an allocator with empty traits + default mem space. + !$omp target uses_allocators(my_alloc) map(tofrom: x, xbuf) defaultmap(none) + !$omp parallel allocate(allocator(my_alloc), align(128): x, xbuf) if(.false.) firstprivate(x, xbuf) + if (mod (TRANSFER (loc(x), iptr), 128) /= 0) & + stop 6 + if (mod (TRANSFER (loc(xbuf), iptr), 128) /= 0) & + stop 7 + !$omp end parallel + !$omp end target + + if (transfer(my_alloc, 0_omp_allocator_handle_kind) /= int(z'ABCD', omp_allocator_handle_kind)) & + stop 8 +end + +! { dg-final { scan-tree-dump-times "#pragma omp target .*private\\(my_alloc\\).*uses_allocators\\(my_alloc: memspace\\(\\), traits\\(trait\\)\\)" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "#pragma omp target .*private\\(my_alloc\\).*uses_allocators\\(my_alloc: memspace\\(\\), traits\\(\\)\\)" 1 "gimple" } }