From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from mail-wr1-x42e.google.com (mail-wr1-x42e.google.com [IPv6:2a00:1450:4864:20::42e]) by sourceware.org (Postfix) with ESMTPS id 09C50385841E for ; Wed, 13 Mar 2024 09:12:27 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 09C50385841E Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=baylibre.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=baylibre.com ARC-Filter: OpenARC Filter v1.0.0 sourceware.org 09C50385841E Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::42e ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1710321153; cv=none; b=HnUNhI2vW2bSwIGhhcyu9jITX1u3Er6b3OvSgYI9UVG12WmP7mrgZIaa+BZ1XJ6NoE7gCgxRST6nKzPuc5r/GYa6+AWhVwG5WaFvTdeLvXYCHCb/50sMI4FCxdWCk3KV5jPjhRV9HtYAahJtfh4oX1J7gWsmk+494IFhjduYjwM= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1710321153; c=relaxed/simple; bh=Skfk+6WvB0XQDwzuUDGcIxfeNKBI5xqZtSPk55ad5ac=; h=DKIM-Signature:From:To:Subject:Date:Message-ID:MIME-Version; b=HmbtJC3P5NrqcOQ19A7TYGk/TyCS2uxiU1yxfU3dbsiZhfIzxfoTXQPJeTR9PxM9c9qvW/rLu7yyLGy+Rxx05g6XrFjBz3doWT96G65GKFsuVHC8o2i4DL+2EkqJvB7tmW3b++7JSo9jSZ807kv3EFHnnmHKBBEosTQnE22mdDo= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-wr1-x42e.google.com with SMTP id ffacd0b85a97d-33d90dfe73cso333698f8f.0 for ; Wed, 13 Mar 2024 02:12:26 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1710321145; x=1710925945; darn=gcc.gnu.org; h=content-transfer-encoding:mime-version:message-id:date:user-agent :references:in-reply-to:subject:cc:to:from:from:to:cc:subject:date :message-id:reply-to; bh=LF8Ef6MUwGH3kvq+KAVVTgY+GyVZWEmb54kHWidHYOo=; b=JtSoY/3okVk25S5fWkcxQNnEkE8ifpQ/XcSS0S0VuMjYH2k1+Ftd117/MkH8Gfu74M jgsA+8AJ0Mah9Kka029V23vwL/ApSjIVV9otGH8ghHW8vUjEkJVCMYhzKnaHF7NK3Eud UK9GU+QbKsSkwLSmHCr9T74+aYdTE3mF8SdhLiMAB2EyfXpSNAb62YyedF57ZhVp3gYK 0AsWjhIWPZht5dhVmxXMVvBxfOnu7aNLcxsy39m8MB+X4RN5nkn7KmR9Z6kTq1U2RIFc vZV6pUmKPmanyp4rkNJNy1unlj9rNnX/3IDDmXGgktgBL+ObLZi/gkk0ZefqsH+umD8N kkKQ== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1710321145; x=1710925945; h=content-transfer-encoding:mime-version:message-id:date:user-agent :references:in-reply-to:subject:cc:to:from:x-gm-message-state:from :to:cc:subject:date:message-id:reply-to; bh=LF8Ef6MUwGH3kvq+KAVVTgY+GyVZWEmb54kHWidHYOo=; b=mG+m0nF1+iIxuwQupXeCAMNWA3rPXy9iEDwPKix7sab49dh5lAYBxY/ZOjXOmEfu1G qXAWc4arNQtMJBYhYY5c1YRTjIBuqop33QsUeyv/2W6MNH9WKCEic6gBPAB0GsyEcl4S 5QK9tzWsTs+gmOIAsY80QJCwf1dcCaI+YT7PvMeKhGbOI3CVIyfT/f5u/R26Gm7d5sMe 3JGC64dpc2LboN6M9Wpk0dt3RQKlxCCGGwLH8E01o9XTswOibFgCoOblBrxoAWYe5lfB igjS/2iNeAys5P8T5PRN2WwZ1RBV5W5dmhR5d2Ew7JbG7iZTydCt220SGdSHsmuZGbT6 NAHw== X-Forwarded-Encrypted: i=1; AJvYcCUim+QPwNVOFTEwYMlj7hv2zcYjElM4II3fN+4WsWxjnD1OTRbi/Rbj9ZW30P6sOoaZQU58iOIt5JPlwkQuSTfE3V5z X-Gm-Message-State: AOJu0YyANKK1F3VutMLaDA3Ac+klEzqTOCF3oFjSv4PPxymRfoZB1lpM vS+TS/ibq6avuMVyPFOW7YkM3glR/JbFop1Zk+eX6EBdMu4fqONIozelWsybIG4= X-Google-Smtp-Source: AGHT+IFURcb74y9AD7MWWxzqK5YOXcyK0jgyetFSS0UpmIKOd7T0ARH7Ujqhnrbr7alRkMD7fUPziQ== X-Received: by 2002:adf:e70d:0:b0:33e:48dd:2927 with SMTP id c13-20020adfe70d000000b0033e48dd2927mr1565280wrm.18.1710321145075; Wed, 13 Mar 2024 02:12:25 -0700 (PDT) Received: from euler.schwinge.homeip.net (p200300c8b70336000b0134869109dcb1.dip0.t-ipconnect.de. [2003:c8:b703:3600:b01:3486:9109:dcb1]) by smtp.gmail.com with ESMTPSA id d9-20020adf9c89000000b0033df46f70dbsm11357301wre.9.2024.03.13.02.12.24 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Wed, 13 Mar 2024 02:12:24 -0700 (PDT) From: Thomas Schwinge To: Chung-Lin Tang Cc: Tobias Burnus , gcc-patches@gcc.gnu.org, fortran@gcc.gnu.org Subject: Re: [PATCH, OpenACC 2.7, v2] readonly modifier support in front-ends In-Reply-To: <2f568f7d-807b-41d1-befb-40039e2edb74@pllab.cs.nthu.edu.tw> References: <87lefaaesb.fsf@euler.schwinge.homeip.net> <87ttqd7m8e.fsf@euler.schwinge.homeip.net> <2f568f7d-807b-41d1-befb-40039e2edb74@pllab.cs.nthu.edu.tw> User-Agent: Notmuch/0.29.3+94~g74c3f1b (https://notmuchmail.org) Emacs/29.1 (x86_64-pc-linux-gnu) Date: Wed, 13 Mar 2024 10:12:17 +0100 Message-ID: <87le6mebri.fsf@euler.schwinge.ddns.net> MIME-Version: 1.0 Content-Type: text/plain; charset=utf-8 Content-Transfer-Encoding: quoted-printable X-Spam-Status: No, score=-10.7 required=5.0 tests=BAYES_00,DKIM_SIGNED,DKIM_VALID,GIT_PATCH_0,RCVD_IN_DNSWL_NONE,SPF_HELO_NONE,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 List-Id: Hi Chung-Lin! On 2024-03-07T17:02:02+0900, Chung-Lin Tang w= rote: > On 2023/10/26 6:43 PM, Thomas Schwinge wrote: >>>>>> +++ b/gcc/tree.h >>>>>> @@ -1813,6 +1813,14 @@ class auto_suppress_location_wrappers >>>>>> #define OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE(NODE) \ >>>>>> (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.addressab= le_flag) >>>>>> >>>>>> +/* Nonzero if OpenACC 'readonly' modifier set, used for 'copyin'. = */ >>>>>> +#define OMP_CLAUSE_MAP_READONLY(NODE) \ >>>>>> + TREE_READONLY (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)) >>>>>> + >>>>>> +/* Same as above, for use in OpenACC cache directives. */ >>>>>> +#define OMP_CLAUSE__CACHE__READONLY(NODE) \ >>>>>> + TREE_READONLY (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE__CACHE_= )) >>>>> I'm not sure if these special accessor functions are actually useful,= or >>>>> we should just directly use 'TREE_READONLY' instead? We're only using >>>>> them in contexts where it's clear that the 'OMP_CLAUSE_SUBCODE_CHECK'= is >>>>> satisfied, for example. >>>> I find directly using TREE_READONLY confusing. >>> >>> FWIW, I've changed to use TREE_NOTHROW instead, if it can give a better= sense of safety :P >>=20 >> I don't understand that, why not use 'TREE_READONLY'? >>=20 >>> I think there's a misunderstanding here anyways: we are not relying on = a DECL marked >>> TREE_READONLY here. We merely need the OMP_CLAUSE_MAP to be marked as O= MP_CLAUSE_MAP_READONLY =3D=3D 1. >>=20 >> Yes, I understand that. My question was why we don't just use >> 'TREE_READONLY (c)', where 'c' is the >> 'OMP_CLAUSE_MAP'/'OMP_CLAUSE__CACHE_' clause (not its decl), and avoid >> the indirection through >> '#define OMP_CLAUSE_MAP_READONLY'/'#define OMP_CLAUSE__CACHE__READONLY', >> given that we're only using them in contexts where it's clear that the >> 'OMP_CLAUSE_SUBCODE_CHECK' is satisfied. I don't have a strong >> preference, though. > > After further re-testing using TREE_NOTHROW, I have reverted to using TRE= E_READONLY ACK, thanks. > because TREE_NOTHROW clashes > with OMP_CLAUSE_RELEASE_DESCRIPTOR (which doesn't use the OMP_CLAUSE_MAP_= * naming convention and is > not documented in gcc/tree-core.h either, hmmm...) Yeah, it's a mess... The same bits of information spread over three different places. (One day I'll turn 'tree's into a proper C++ class hierarchy, with accessor methods for such flags, statically checked at compile-time, and thus documented in a single place. Etc.) > I have added the comment adjustments in gcc/tree-core.h for the new uses = of TREE_READONLY/readonly_flag. > > We basically all use OMP_CLAUSE_SUBCODE_CHECK macros for OpenMP clause ex= pressions exclusively, > so I don't see a reason to diverge from that style (even when context is = clear). ACK. > I have greatly expanded the test scan patterns to include parallel/kernel= s/serial/data/enter data, > as well as non-readonly copyin clause together with readonly. Thanks. > Also added simple 'declare' tests, but there is not anything to scan in t= he 'tree-original' dump though. Yeah, the current OpenACC 'declare' implementation is "special". >>> --- a/gcc/fortran/openmp.cc >>> +++ b/gcc/fortran/openmp.cc >>> @@ -1197,7 +1197,7 @@ omp_inv_mask::omp_inv_mask (const omp_mask &m) : = omp_mask (m) >>> >>> static bool >>> gfc_match_omp_map_clause (gfc_omp_namelist **list, gfc_omp_map_op map_= op, >>> - bool allow_common, bool allow_derived) >>> + bool allow_common, bool allow_derived, bool rea= donly =3D false) >>> { >>> gfc_omp_namelist **head =3D NULL; >>> if (gfc_match_omp_variable_list ("", list, allow_common, NULL, &head= , true, >>> @@ -1206,7 +1206,10 @@ gfc_match_omp_map_clause (gfc_omp_namelist **lis= t, gfc_omp_map_op map_op, >>> { >>> gfc_omp_namelist *n; >>> for (n =3D *head; n; n =3D n->next) >>> - n->u.map_op =3D map_op; >>> + { >>> + n->u.map.op =3D map_op; >>> + n->u.map.readonly =3D readonly; >>> + } >>> return true; >>> } >>=20 >> Didn't we conclude that "not doing it here is cleaner" (Tobias' words), >> and instead do this "Similar to 'c_parser_omp_var_list_parens'" (my >> words)? That is, not add the 'bool readonly' formal parameter to >> 'gfc_match_omp_map_clause'. > > Fixed in this v3 patch. Thanks. > Again, tested on x86_64-linux + nvptx offloading. Okay for mainline? Yes, thanks. Gr=C3=BC=C3=9Fe Thomas > gcc/c/ChangeLog: > > * c-parser.cc (c_parser_oacc_data_clause): Add parsing support for > 'readonly' modifier, set OMP_CLAUSE_MAP_READONLY if readonly modifier > found, update comments. > (c_parser_oacc_cache): Add parsing support for 'readonly' modifier, > set OMP_CLAUSE__CACHE__READONLY if readonly modifier found, update > comments. > > gcc/cp/ChangeLog: > > * parser.cc (cp_parser_oacc_data_clause): Add parsing support for > 'readonly' modifier, set OMP_CLAUSE_MAP_READONLY if readonly modifier > found, update comments. > (cp_parser_oacc_cache): Add parsing support for 'readonly' modifier, > set OMP_CLAUSE__CACHE__READONLY if readonly modifier found, update > comments. > > gcc/fortran/ChangeLog: > > * dump-parse-tree.cc (show_omp_namelist): Print "readonly," for > OMP_LIST_MAP and OMP_LIST_CACHE if n->u.map.readonly is set. > Adjust 'n->u.map_op' to 'n->u.map.op'. > * gfortran.h (typedef struct gfc_omp_namelist): Adjust map_op as > 'ENUM_BITFIELD (gfc_omp_map_op) op:8', add 'bool readonly' field, > change to named struct field 'map'. > > * openmp.cc (gfc_match_omp_map_clause): Adjust 'n->u.map_op' to > 'n->u.map.op'. > (gfc_match_omp_clause_reduction): Likewise. > > (gfc_match_omp_clauses): Add readonly modifier parsing for OpenACC > copyin clause, set 'n->u.map.op' and 'n->u.map.readonly' for parsed > clause. Adjust 'n->u.map_op' to 'n->u.map.op'. > (gfc_match_oacc_declare): Adjust 'n->u.map_op' to 'n->u.map.op'. > (gfc_match_oacc_cache): Add readonly modifier parsing for OpenACC > cache directive. > (resolve_omp_clauses): Adjust 'n->u.map_op' to 'n->u.map.op'. > * trans-decl.cc (add_clause): Adjust 'n->u.map_op' to 'n->u.map.op'. > (finish_oacc_declare): Likewise. > * trans-openmp.cc (gfc_trans_omp_clauses): Set OMP_CLAUSE_MAP_READONLY, > OMP_CLAUSE__CACHE__READONLY to 1 when readonly is set. Adjust > 'n->u.map_op' to 'n->u.map.op'. > (gfc_add_clause_implicitly): Adjust 'n->u.map_op' to 'n->u.map.op'. > > gcc/ChangeLog: > * tree.h (OMP_CLAUSE_MAP_READONLY): New macro. > (OMP_CLAUSE__CACHE__READONLY): New macro. > * tree-core.h (struct GTY(()) tree_base): Adjust comments for new > uses of readonly_flag bit in OMP_CLAUSE_MAP_READONLY and > OMP_CLAUSE__CACHE__READONLY. > * tree-pretty-print.cc (dump_omp_clause): Add support for printing > OMP_CLAUSE_MAP_READONLY and OMP_CLAUSE__CACHE__READONLY. > > gcc/testsuite/ChangeLog: > > * c-c++-common/goacc/readonly-1.c: New test. > * gfortran.dg/goacc/readonly-1.f90: New test. > > > > > > diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc > index 53e99aa29d9..00f8bf4376e 100644 > --- a/gcc/c/c-parser.cc > +++ b/gcc/c/c-parser.cc > @@ -15627,7 +15627,11 @@ c_parser_omp_var_list_parens (c_parser *parser, = enum omp_clause_code kind, > OpenACC 2.6: > no_create ( variable-list ) > attach ( variable-list ) > - detach ( variable-list ) */ > + detach ( variable-list ) > + > + OpenACC 2.7: > + copyin (readonly : variable-list ) > + */ >=20=20 > static tree > c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind, > @@ -15680,11 +15684,37 @@ c_parser_oacc_data_clause (c_parser *parser, pr= agma_omp_clause c_kind, > default: > gcc_unreachable (); > } > - tree nl, c; > - nl =3D c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list, fal= se); >=20=20 > - for (c =3D nl; c !=3D list; c =3D OMP_CLAUSE_CHAIN (c)) > - OMP_CLAUSE_SET_MAP_KIND (c, kind); > + tree nl =3D list; > + bool readonly =3D false; > + location_t open_loc =3D c_parser_peek_token (parser)->location; > + matching_parens parens; > + if (parens.require_open (parser)) > + { > + /* Turn on readonly modifier parsing for copyin clause. */ > + if (c_kind =3D=3D PRAGMA_OACC_CLAUSE_COPYIN) > + { > + c_token *token =3D c_parser_peek_token (parser); > + if (token->type =3D=3D CPP_NAME > + && !strcmp (IDENTIFIER_POINTER (token->value), "readonly") > + && c_parser_peek_2nd_token (parser)->type =3D=3D CPP_COLON) > + { > + c_parser_consume_token (parser); > + c_parser_consume_token (parser); > + readonly =3D true; > + } > + } > + nl =3D c_parser_omp_variable_list (parser, open_loc, OMP_CLAUSE_MA= P, list, > + false); > + parens.skip_until_found_close (parser); > + } > + > + for (tree c =3D nl; c !=3D list; c =3D OMP_CLAUSE_CHAIN (c)) > + { > + OMP_CLAUSE_SET_MAP_KIND (c, kind); > + if (readonly) > + OMP_CLAUSE_MAP_READONLY (c) =3D 1; > + } >=20=20 > return nl; > } > @@ -19821,15 +19851,39 @@ c_parser_omp_structured_block (c_parser *parser= , bool *if_p) > /* OpenACC 2.0: > # pragma acc cache (variable-list) new-line >=20=20 > + OpenACC 2.7: > + # pragma acc cache (readonly: variable-list) new-line > + > LOC is the location of the #pragma token. > */ >=20=20 > static tree > c_parser_oacc_cache (location_t loc, c_parser *parser) > { > - tree stmt, clauses; > + tree stmt, clauses =3D NULL_TREE; > + bool readonly =3D false; > + location_t open_loc =3D c_parser_peek_token (parser)->location; > + matching_parens parens; > + if (parens.require_open (parser)) > + { > + c_token *token =3D c_parser_peek_token (parser); > + if (token->type =3D=3D CPP_NAME > + && !strcmp (IDENTIFIER_POINTER (token->value), "readonly") > + && c_parser_peek_2nd_token (parser)->type =3D=3D CPP_COLON) > + { > + c_parser_consume_token (parser); > + c_parser_consume_token (parser); > + readonly =3D true; > + } > + clauses =3D c_parser_omp_variable_list (parser, open_loc, > + OMP_CLAUSE__CACHE_, NULL_TREE); > + parens.skip_until_found_close (parser); > + } > + > + if (readonly) > + for (tree c =3D clauses; c; c =3D OMP_CLAUSE_CHAIN (c)) > + OMP_CLAUSE__CACHE__READONLY (c) =3D 1; >=20=20 > - clauses =3D c_parser_omp_var_list_parens (parser, OMP_CLAUSE__CACHE_, = NULL); > clauses =3D c_finish_omp_clauses (clauses, C_ORT_ACC); >=20=20 > c_parser_skip_to_pragma_eol (parser); > diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc > index e32acfc30a2..4fe27fb07b2 100644 > --- a/gcc/cp/parser.cc > +++ b/gcc/cp/parser.cc > @@ -38544,7 +38544,11 @@ cp_parser_omp_var_list (cp_parser *parser, enum = omp_clause_code kind, tree list, > OpenACC 2.6: > no_create ( variable-list ) > attach ( variable-list ) > - detach ( variable-list ) */ > + detach ( variable-list ) > + > + OpenACC 2.7: > + copyin (readonly : variable-list ) > + */ >=20=20 > static tree > cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind, > @@ -38597,11 +38601,34 @@ cp_parser_oacc_data_clause (cp_parser *parser, = pragma_omp_clause c_kind, > default: > gcc_unreachable (); > } > - tree nl, c; > - nl =3D cp_parser_omp_var_list (parser, OMP_CLAUSE_MAP, list, false); >=20=20 > - for (c =3D nl; c !=3D list; c =3D OMP_CLAUSE_CHAIN (c)) > - OMP_CLAUSE_SET_MAP_KIND (c, kind); > + tree nl =3D list; > + bool readonly =3D false; > + if (cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN)) > + { > + /* Turn on readonly modifier parsing for copyin clause. */ > + if (c_kind =3D=3D PRAGMA_OACC_CLAUSE_COPYIN) > + { > + cp_token *token =3D cp_lexer_peek_token (parser->lexer); > + if (token->type =3D=3D CPP_NAME > + && !strcmp (IDENTIFIER_POINTER (token->u.value), "readonly") > + && cp_lexer_peek_nth_token (parser->lexer, 2)->type =3D=3D CPP_CO= LON) > + { > + cp_lexer_consume_token (parser->lexer); > + cp_lexer_consume_token (parser->lexer); > + readonly =3D true; > + } > + } > + nl =3D cp_parser_omp_var_list_no_open (parser, OMP_CLAUSE_MAP, lis= t, NULL, > + false); > + } > + > + for (tree c =3D nl; c !=3D list; c =3D OMP_CLAUSE_CHAIN (c)) > + { > + OMP_CLAUSE_SET_MAP_KIND (c, kind); > + if (readonly) > + OMP_CLAUSE_MAP_READONLY (c) =3D 1; > + } >=20=20 > return nl; > } > @@ -47178,6 +47205,9 @@ cp_parser_omp_target (cp_parser *parser, cp_token= *pragma_tok, >=20=20 > /* OpenACC 2.0: > # pragma acc cache (variable-list) new-line > + > + OpenACC 2.7: > + # pragma acc cache (readonly: variable-list) new-line > */ >=20=20 > static tree > @@ -47187,9 +47217,28 @@ cp_parser_oacc_cache (cp_parser *parser, cp_toke= n *pragma_tok) > clauses. */ > auto_suppress_location_wrappers sentinel; >=20=20 > - tree stmt, clauses; > + tree stmt, clauses =3D NULL_TREE; > + bool readonly =3D false; > + > + if (cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN)) > + { > + cp_token *token =3D cp_lexer_peek_token (parser->lexer); > + if (token->type =3D=3D CPP_NAME > + && !strcmp (IDENTIFIER_POINTER (token->u.value), "readonly") > + && cp_lexer_peek_nth_token (parser->lexer, 2)->type =3D=3D CPP_COLON) > + { > + cp_lexer_consume_token (parser->lexer); > + cp_lexer_consume_token (parser->lexer); > + readonly =3D true; > + } > + clauses =3D cp_parser_omp_var_list_no_open (parser, OMP_CLAUSE__CA= CHE_, > + NULL, NULL); > + } > + > + if (readonly) > + for (tree c =3D clauses; c; c =3D OMP_CLAUSE_CHAIN (c)) > + OMP_CLAUSE__CACHE__READONLY (c) =3D 1; >=20=20 > - clauses =3D cp_parser_omp_var_list (parser, OMP_CLAUSE__CACHE_, NULL_T= REE); > clauses =3D finish_omp_clauses (clauses, C_ORT_ACC); >=20=20 > cp_parser_require_pragma_eol (parser, cp_lexer_peek_token (parser->lex= er)); > diff --git a/gcc/fortran/dump-parse-tree.cc b/gcc/fortran/dump-parse-tree= .cc > index 7b154eb3ca7..db84b06289b 100644 > --- a/gcc/fortran/dump-parse-tree.cc > +++ b/gcc/fortran/dump-parse-tree.cc > @@ -1400,6 +1400,9 @@ show_omp_namelist (int list_type, gfc_omp_namelist = *n) > fputs (") ALLOCATE(", dumpfile); > continue; > } > + if ((list_type =3D=3D OMP_LIST_MAP || list_type =3D=3D OMP_LIST_CA= CHE) > + && n->u.map.readonly) > + fputs ("readonly,", dumpfile); > if (list_type =3D=3D OMP_LIST_REDUCTION) > switch (n->u.reduction_op) > { > @@ -1467,7 +1470,7 @@ show_omp_namelist (int list_type, gfc_omp_namelist = *n) > default: break; > } > else if (list_type =3D=3D OMP_LIST_MAP) > - switch (n->u.map_op) > + switch (n->u.map.op) > { > case OMP_MAP_ALLOC: fputs ("alloc:", dumpfile); break; > case OMP_MAP_TO: fputs ("to:", dumpfile); break; > diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h > index ebba2336e12..32b792f85fb 100644 > --- a/gcc/fortran/gfortran.h > +++ b/gcc/fortran/gfortran.h > @@ -1363,7 +1363,11 @@ typedef struct gfc_omp_namelist > { > gfc_omp_reduction_op reduction_op; > gfc_omp_depend_doacross_op depend_doacross_op; > - gfc_omp_map_op map_op; > + struct > + { > + ENUM_BITFIELD (gfc_omp_map_op) op:8; > + bool readonly; > + } map; > gfc_expr *align; > struct > { > diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc > index 38de60238c0..5c44e666eb9 100644 > --- a/gcc/fortran/openmp.cc > +++ b/gcc/fortran/openmp.cc > @@ -1210,7 +1210,7 @@ gfc_match_omp_map_clause (gfc_omp_namelist **list, = gfc_omp_map_op map_op, > { > gfc_omp_namelist *n; > for (n =3D *head; n; n =3D n->next) > - n->u.map_op =3D map_op; > + n->u.map.op =3D map_op; > return true; > } >=20=20 > @@ -1524,7 +1524,7 @@ gfc_match_omp_clause_reduction (char pc, gfc_omp_cl= auses *c, bool openacc, > gfc_omp_namelist *p =3D gfc_get_omp_namelist (), **tl; > p->sym =3D n->sym; > p->where =3D p->where; > - p->u.map_op =3D OMP_MAP_ALWAYS_TOFROM; > + p->u.map.op =3D OMP_MAP_ALWAYS_TOFROM; >=20=20 > tl =3D &c->lists[OMP_LIST_MAP]; > while (*tl) > @@ -2181,11 +2181,25 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, cons= t omp_mask mask, > { > if (openacc) > { > - if (gfc_match ("copyin ( ") =3D=3D MATCH_YES > - && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], > - OMP_MAP_TO, true, > - allow_derived)) > - continue; > + if (gfc_match ("copyin ( ") =3D=3D MATCH_YES) > + { > + bool readonly =3D gfc_match ("readonly : ") =3D=3D MATCH_YES; > + head =3D NULL; > + if (gfc_match_omp_variable_list ("", > + &c->lists[OMP_LIST_MAP], > + true, NULL, &head, true, > + allow_derived) > + =3D=3D MATCH_YES) > + { > + gfc_omp_namelist *n; > + for (n =3D *head; n; n =3D n->next) > + { > + n->u.map.op =3D OMP_MAP_TO; > + n->u.map.readonly =3D readonly; > + } > + continue; > + } > + } > } > else if (gfc_match_omp_variable_list ("copyin (", > &c->lists[OMP_LIST_COPYIN], > @@ -3134,7 +3148,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const = omp_mask mask, > { > gfc_omp_namelist *n; > for (n =3D *head; n; n =3D n->next) > - n->u.map_op =3D map_op; > + n->u.map.op =3D map_op; > continue; > } > gfc_current_locus =3D old_loc; > @@ -4002,7 +4016,7 @@ gfc_match_oacc_declare (void) > if (gfc_current_ns->proc_name > && gfc_current_ns->proc_name->attr.flavor =3D=3D FL_MODULE) > { > - if (n->u.map_op !=3D OMP_MAP_ALLOC && n->u.map_op !=3D OMP_MAP_TO) > + if (n->u.map.op !=3D OMP_MAP_ALLOC && n->u.map.op !=3D OMP_MAP_TO) > { > gfc_error ("Invalid clause in module with !$ACC DECLARE at %L", > &where); > @@ -4036,7 +4050,7 @@ gfc_match_oacc_declare (void) > return MATCH_ERROR; > } >=20=20 > - switch (n->u.map_op) > + switch (n->u.map.op) > { > case OMP_MAP_FORCE_ALLOC: > case OMP_MAP_ALLOC: > @@ -4151,21 +4165,36 @@ gfc_match_oacc_wait (void) > match > gfc_match_oacc_cache (void) > { > + bool readonly =3D false; > gfc_omp_clauses *c =3D gfc_get_omp_clauses (); > /* The OpenACC cache directive explicitly only allows "array elements = or > subarrays", which we're currently not checking here. Either check = this > after the call of gfc_match_omp_variable_list, or add something lik= e a > only_sections variant next to its allow_sections parameter. */ > - match m =3D gfc_match_omp_variable_list (" (", > - &c->lists[OMP_LIST_CACHE], true, > - NULL, NULL, true); > + match m =3D gfc_match (" ( "); > if (m !=3D MATCH_YES) > { > gfc_free_omp_clauses(c); > return m; > } >=20=20 > - if (gfc_current_state() !=3D COMP_DO=20 > + if (gfc_match ("readonly : ") =3D=3D MATCH_YES) > + readonly =3D true; > + > + gfc_omp_namelist **head =3D NULL; > + m =3D gfc_match_omp_variable_list ("", &c->lists[OMP_LIST_CACHE], true, > + NULL, &head, true); > + if (m !=3D MATCH_YES) > + { > + gfc_free_omp_clauses(c); > + return m; > + } > + > + if (readonly) > + for (gfc_omp_namelist *n =3D *head; n; n =3D n->next) > + n->u.map.readonly =3D true; > + > + if (gfc_current_state() !=3D COMP_DO > && gfc_current_state() !=3D COMP_DO_CONCURRENT) > { > gfc_error ("ACC CACHE directive must be inside of loop %C"); > @@ -8436,8 +8465,8 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clause= s *omp_clauses, > } > if (openacc > && list =3D=3D OMP_LIST_MAP > - && (n->u.map_op =3D=3D OMP_MAP_ATTACH > - || n->u.map_op =3D=3D OMP_MAP_DETACH)) > + && (n->u.map.op =3D=3D OMP_MAP_ATTACH > + || n->u.map.op =3D=3D OMP_MAP_DETACH)) > { > symbol_attribute attr; > if (n->expr) > @@ -8447,7 +8476,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clause= s *omp_clauses, > if (!attr.pointer && !attr.allocatable) > gfc_error ("%qs clause argument must be ALLOCATABLE or " > "a POINTER at %L", > - (n->u.map_op =3D=3D OMP_MAP_ATTACH) ? "attach" > + (n->u.map.op =3D=3D OMP_MAP_ATTACH) ? "attach" > : "detach", &n->where); > } > if (lastref > @@ -8518,7 +8547,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clause= s *omp_clauses, > else if (openacc) > { > if (list =3D=3D OMP_LIST_MAP > - && n->u.map_op =3D=3D OMP_MAP_FORCE_DEVICEPTR) > + && n->u.map.op =3D=3D OMP_MAP_FORCE_DEVICEPTR) > resolve_oacc_deviceptr_clause (n->sym, n->where, name); > else > resolve_oacc_data_clauses (n->sym, n->where, name); > @@ -8540,7 +8569,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clause= s *omp_clauses, > { > case EXEC_OMP_TARGET: > case EXEC_OMP_TARGET_DATA: > - switch (n->u.map_op) > + switch (n->u.map.op) > { > case OMP_MAP_TO: > case OMP_MAP_ALWAYS_TO: > @@ -8567,7 +8596,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clause= s *omp_clauses, > } > break; > case EXEC_OMP_TARGET_ENTER_DATA: > - switch (n->u.map_op) > + switch (n->u.map.op) > { > case OMP_MAP_TO: > case OMP_MAP_ALWAYS_TO: > @@ -8577,16 +8606,16 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clau= ses *omp_clauses, > case OMP_MAP_PRESENT_ALLOC: > break; > case OMP_MAP_TOFROM: > - n->u.map_op =3D OMP_MAP_TO; > + n->u.map.op =3D OMP_MAP_TO; > break; > case OMP_MAP_ALWAYS_TOFROM: > - n->u.map_op =3D OMP_MAP_ALWAYS_TO; > + n->u.map.op =3D OMP_MAP_ALWAYS_TO; > break; > case OMP_MAP_PRESENT_TOFROM: > - n->u.map_op =3D OMP_MAP_PRESENT_TO; > + n->u.map.op =3D OMP_MAP_PRESENT_TO; > break; > case OMP_MAP_ALWAYS_PRESENT_TOFROM: > - n->u.map_op =3D OMP_MAP_ALWAYS_PRESENT_TO; > + n->u.map.op =3D OMP_MAP_ALWAYS_PRESENT_TO; > break; > default: > gfc_error ("TARGET ENTER DATA with map-type other " > @@ -8596,7 +8625,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clause= s *omp_clauses, > } > break; > case EXEC_OMP_TARGET_EXIT_DATA: > - switch (n->u.map_op) > + switch (n->u.map.op) > { > case OMP_MAP_FROM: > case OMP_MAP_ALWAYS_FROM: > @@ -8606,16 +8635,16 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clau= ses *omp_clauses, > case OMP_MAP_DELETE: > break; > case OMP_MAP_TOFROM: > - n->u.map_op =3D OMP_MAP_FROM; > + n->u.map.op =3D OMP_MAP_FROM; > break; > case OMP_MAP_ALWAYS_TOFROM: > - n->u.map_op =3D OMP_MAP_ALWAYS_FROM; > + n->u.map.op =3D OMP_MAP_ALWAYS_FROM; > break; > case OMP_MAP_PRESENT_TOFROM: > - n->u.map_op =3D OMP_MAP_PRESENT_FROM; > + n->u.map.op =3D OMP_MAP_PRESENT_FROM; > break; > case OMP_MAP_ALWAYS_PRESENT_TOFROM: > - n->u.map_op =3D OMP_MAP_ALWAYS_PRESENT_FROM; > + n->u.map.op =3D OMP_MAP_ALWAYS_PRESENT_FROM; > break; > default: > gfc_error ("TARGET EXIT DATA with map-type other " > diff --git a/gcc/fortran/trans-decl.cc b/gcc/fortran/trans-decl.cc > index 6d463036966..b7dea11461f 100644 > --- a/gcc/fortran/trans-decl.cc > +++ b/gcc/fortran/trans-decl.cc > @@ -6744,7 +6744,7 @@ add_clause (gfc_symbol *sym, gfc_omp_map_op map_op) >=20=20 > n =3D gfc_get_omp_namelist (); > n->sym =3D sym; > - n->u.map_op =3D map_op; > + n->u.map.op =3D map_op; >=20=20 > if (!module_oacc_clauses) > module_oacc_clauses =3D gfc_get_omp_clauses (); > @@ -6846,10 +6846,10 @@ finish_oacc_declare (gfc_namespace *ns, gfc_symbo= l *sym, bool block) >=20=20 > for (n =3D omp_clauses->lists[OMP_LIST_MAP]; n; n =3D n->next) > { > - switch (n->u.map_op) > + switch (n->u.map.op) > { > case OMP_MAP_DEVICE_RESIDENT: > - n->u.map_op =3D OMP_MAP_FORCE_ALLOC; > + n->u.map.op =3D OMP_MAP_FORCE_ALLOC; > break; >=20=20 > default: > diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc > index a2bf15665b3..fa1bfd41380 100644 > --- a/gcc/fortran/trans-openmp.cc > +++ b/gcc/fortran/trans-openmp.cc > @@ -3139,7 +3139,10 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp= _clauses *clauses, > || (n->expr && gfc_expr_attr (n->expr).pointer))) > always_modifier =3D true; >=20=20 > - switch (n->u.map_op) > + if (n->u.map.readonly) > + OMP_CLAUSE_MAP_READONLY (node) =3D 1; > + > + switch (n->u.map.op) > { > case OMP_MAP_ALLOC: > OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALLOC); > @@ -3266,8 +3269,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_= clauses *clauses, > && n->sym->attr.omp_declare_target > && (always_modifier || n->sym->attr.pointer) > && op !=3D EXEC_OMP_TARGET_EXIT_DATA > - && n->u.map_op !=3D OMP_MAP_DELETE > - && n->u.map_op !=3D OMP_MAP_RELEASE) > + && n->u.map.op !=3D OMP_MAP_DELETE > + && n->u.map.op !=3D OMP_MAP_RELEASE) > { > gcc_assert (n->sym->ts.u.cl->backend_decl); > node5 =3D build_omp_clause (input_location, OMP_CLAUSE_MAP); > @@ -3333,7 +3336,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_= clauses *clauses, > { > enum gomp_map_kind gmk =3D GOMP_MAP_POINTER; > if (op =3D=3D EXEC_OMP_TARGET_EXIT_DATA > - && n->u.map_op =3D=3D OMP_MAP_DELETE) > + && n->u.map.op =3D=3D OMP_MAP_DELETE) > gmk =3D GOMP_MAP_DELETE; > else if (op =3D=3D EXEC_OMP_TARGET_EXIT_DATA) > gmk =3D GOMP_MAP_RELEASE; > @@ -3356,7 +3359,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_= clauses *clauses, > { > enum gomp_map_kind gmk; > if (op =3D=3D EXEC_OMP_TARGET_EXIT_DATA > - && n->u.map_op =3D=3D OMP_MAP_DELETE) > + && n->u.map.op =3D=3D OMP_MAP_DELETE) > gmk =3D GOMP_MAP_DELETE; > else if (op =3D=3D EXEC_OMP_TARGET_EXIT_DATA) > gmk =3D GOMP_MAP_RELEASE; > @@ -3388,18 +3391,18 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_om= p_clauses *clauses, > node2 =3D build_omp_clause (input_location, OMP_CLAUSE_MAP); > OMP_CLAUSE_DECL (node2) =3D decl; > OMP_CLAUSE_SIZE (node2) =3D TYPE_SIZE_UNIT (type); > - if (n->u.map_op =3D=3D OMP_MAP_DELETE) > + if (n->u.map.op =3D=3D OMP_MAP_DELETE) > map_kind =3D GOMP_MAP_DELETE; > else if (op =3D=3D EXEC_OMP_TARGET_EXIT_DATA > - || n->u.map_op =3D=3D OMP_MAP_RELEASE) > + || n->u.map.op =3D=3D OMP_MAP_RELEASE) > map_kind =3D GOMP_MAP_RELEASE; > else > map_kind =3D GOMP_MAP_TO_PSET; > OMP_CLAUSE_SET_MAP_KIND (node2, map_kind); >=20=20 > if (op !=3D EXEC_OMP_TARGET_EXIT_DATA > - && n->u.map_op !=3D OMP_MAP_DELETE > - && n->u.map_op !=3D OMP_MAP_RELEASE) > + && n->u.map.op !=3D OMP_MAP_DELETE > + && n->u.map.op !=3D OMP_MAP_RELEASE) > { > node3 =3D build_omp_clause (input_location, > OMP_CLAUSE_MAP); > @@ -3417,7 +3420,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_= clauses *clauses, > =3D gfc_conv_descriptor_data_get (decl); > OMP_CLAUSE_SIZE (node3) =3D size_int (0); >=20=20 > - if (n->u.map_op =3D=3D OMP_MAP_ATTACH) > + if (n->u.map.op =3D=3D OMP_MAP_ATTACH) > { > /* Standalone attach clauses used with arrays with > descriptors must copy the descriptor to the > @@ -3433,7 +3436,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_= clauses *clauses, > node3 =3D NULL; > goto finalize_map_clause; > } > - else if (n->u.map_op =3D=3D OMP_MAP_DETACH) > + else if (n->u.map.op =3D=3D OMP_MAP_DETACH) > { > OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_DETACH); > /* Similarly to above, we don't want to unmap PTR > @@ -3626,8 +3629,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_= clauses *clauses, > to perform a single attach/detach operation, of the > pointer itself, not of the pointed-to object. */ > if (openacc > - && (n->u.map_op =3D=3D OMP_MAP_ATTACH > - || n->u.map_op =3D=3D OMP_MAP_DETACH)) > + && (n->u.map.op =3D=3D OMP_MAP_ATTACH > + || n->u.map.op =3D=3D OMP_MAP_DETACH)) > { > OMP_CLAUSE_DECL (node) > =3D build_fold_addr_expr (OMP_CLAUSE_DECL (node)); > @@ -3656,7 +3659,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_= clauses *clauses, > se.string_length), > TYPE_SIZE_UNIT (tmp)); > gomp_map_kind kind; > - if (n->u.map_op =3D=3D OMP_MAP_DELETE) > + if (n->u.map.op =3D=3D OMP_MAP_DELETE) > kind =3D GOMP_MAP_DELETE; > else if (op =3D=3D EXEC_OMP_TARGET_EXIT_DATA) > kind =3D GOMP_MAP_RELEASE; > @@ -3713,8 +3716,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_= clauses *clauses, > to perform a single attach/detach operation, of the > pointer itself, not of the pointed-to object. */ > if (openacc > - && (n->u.map_op =3D=3D OMP_MAP_ATTACH > - || n->u.map_op =3D=3D OMP_MAP_DETACH)) > + && (n->u.map.op =3D=3D OMP_MAP_ATTACH > + || n->u.map.op =3D=3D OMP_MAP_DETACH)) > { > OMP_CLAUSE_DECL (node) > =3D build_fold_addr_expr (inner); > @@ -3806,8 +3809,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_= clauses *clauses, > { > /* Bare attach and detach clauses don't want any > additional nodes. */ > - if ((n->u.map_op =3D=3D OMP_MAP_ATTACH > - || n->u.map_op =3D=3D OMP_MAP_DETACH) > + if ((n->u.map.op =3D=3D OMP_MAP_ATTACH > + || n->u.map.op =3D=3D OMP_MAP_DETACH) > && (POINTER_TYPE_P (TREE_TYPE (inner)) > || GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (inner)))) > { > @@ -3840,8 +3843,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_= clauses *clauses, > map_kind =3D ((GOMP_MAP_ALWAYS_P (map_kind) > || gfc_expr_attr (n->expr).pointer) > ? GOMP_MAP_ALWAYS_TO : GOMP_MAP_TO); > - else if (n->u.map_op =3D=3D OMP_MAP_RELEASE > - || n->u.map_op =3D=3D OMP_MAP_DELETE) > + else if (n->u.map.op =3D=3D OMP_MAP_RELEASE > + || n->u.map.op =3D=3D OMP_MAP_DELETE) > ; > else if (op =3D=3D EXEC_OMP_TARGET_EXIT_DATA > || op =3D=3D EXEC_OACC_EXIT_DATA) > @@ -4088,6 +4091,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_= clauses *clauses, > } > if (n->u.present_modifier) > OMP_CLAUSE_MOTION_PRESENT (node) =3D 1; > + if (list =3D=3D OMP_LIST_CACHE && n->u.map.readonly) > + OMP_CLAUSE__CACHE__READONLY (node) =3D 1; > omp_clauses =3D gfc_trans_add_clause (node, omp_clauses); > } > break; > @@ -6561,7 +6566,7 @@ gfc_add_clause_implicitly (gfc_omp_clauses *clauses= _out, > n2->where =3D n->where; > n2->sym =3D n->sym; > if (is_target) > - n2->u.map_op =3D OMP_MAP_TOFROM; > + n2->u.map.op =3D OMP_MAP_TOFROM; > if (tail) > { > tail->next =3D n2; > diff --git a/gcc/testsuite/c-c++-common/goacc/readonly-1.c b/gcc/testsuit= e/c-c++-common/goacc/readonly-1.c > new file mode 100644 > index 00000000000..34fc92c24d5 > --- /dev/null > +++ b/gcc/testsuite/c-c++-common/goacc/readonly-1.c > @@ -0,0 +1,59 @@ > +/* { dg-additional-options "-fdump-tree-original" } */ > + > +struct S > +{ > + int *ptr; > + float f; > +}; > + > +int a[32], b[32]; > +#pragma acc declare copyin(readonly: a) copyin(b) > + > +int main (void) > +{ > + int x[32], y[32]; > + struct S s =3D {x, 0}; > + > + #pragma acc parallel copyin(readonly: x[:32], s.ptr[:16]) copyin(y[:32= ]) > + { > + #pragma acc cache (readonly: x[:32]) > + #pragma acc cache (y[:32]) > + } > + > + #pragma acc kernels copyin(readonly: x[:32], s.ptr[:16]) copyin(y[:32]) > + { > + #pragma acc cache (readonly: x[:32]) > + #pragma acc cache (y[:32]) > + } > + > + #pragma acc serial copyin(readonly: x[:32], s.ptr[:16]) copyin(y[:32]) > + { > + #pragma acc cache (readonly: x[:32]) > + #pragma acc cache (y[:32]) > + } > + > + #pragma acc data copyin(readonly: x[:32], s.ptr[:16]) copyin(y[:32]) > + { > + #pragma acc cache (readonly: x[:32]) > + #pragma acc cache (y[:32]) > + } > + > + #pragma acc enter data copyin(readonly: x[:32], s.ptr[:16]) copyin(y[:= 32]) > + > + return 0; > +} > + > +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(to:= y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: = \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1= "original" { target { c } } } } */ > +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc kernels map\\(to:y= \\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \= [0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 = "original" { target { c } } } } */ > +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc serial map\\(to:y\= \\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*s.ptr \\\[len: \[= 0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "= original" { target { c } } } } */ > +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(to:y\\\= [0\\\] \\\[len: \[0-9\]+\\\]\\) map\\(readonly,to:\\*s.ptr \\\[len: \[0-9\]= +\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 "origi= nal" { target { c } } } } */ > +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(t= o:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) map\\(readonly,to:\\*s.ptr \\\[len: \= [0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\)" 1 = "original" { target { c } } } } */ > + > +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(to:= y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*NON_LVALUE_EXPR= \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: = \[0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */ > +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc kernels map\\(to:y= \\\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*NON_LVALUE_EXPR = \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \= [0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */ > +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc serial map\\(to:y\= \\[0\\\] \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:\\*NON_LVALUE_EXPR <= s.ptr> \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[= 0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */ > +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(to:y\\\= [0\\\] \\\[len: \[0-9\]+\\\]\\) map\\(readonly,to:\\*NON_LVALUE_EXPR \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \[0-9\]= +\\\]\\)" 1 "original" { target { c++ } } } } */ > +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(t= o:y\\\[0\\\] \\\[len: \[0-9\]+\\\]\\) map\\(readonly,to:\\*NON_LVALUE_EXPR = \\\[len: \[0-9\]+\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: \= [0-9\]+\\\]\\)" 1 "original" { target { c++ } } } } */ > + > +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(readonly:= x\\\[0\\\] \\\[len: \[0-9\]+\\\]\\);$" 4 "original" } } */ > +/* { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(y\\\[0\\\= ] \\\[len: \[0-9\]+\\\]\\);$" 4 "original" } } */ > diff --git a/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90 b/gcc/testsui= te/gfortran.dg/goacc/readonly-1.f90 > new file mode 100644 > index 00000000000..696ebd08321 > --- /dev/null > +++ b/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90 > @@ -0,0 +1,89 @@ > +! { dg-additional-options "-fdump-tree-original" } > + > +subroutine foo (a, n) > + integer :: n, a(:) > + integer :: i, b(n), c(n) > + !$acc parallel copyin(readonly: a(:), b(:n)) copyin(c(:)) > + do i =3D 1,32 > + !$acc cache (readonly: a(:), b(:n)) > + !$acc cache (c(:)) > + enddo > + !$acc end parallel > + > + !$acc kernels copyin(readonly: a(:), b(:n)) copyin(c(:)) > + do i =3D 1,32 > + !$acc cache (readonly: a(:), b(:n)) > + !$acc cache (c(:)) > + enddo > + !$acc end kernels > + > + !$acc serial copyin(readonly: a(:), b(:n)) copyin(c(:)) > + do i =3D 1,32 > + !$acc cache (readonly: a(:), b(:n)) > + !$acc cache (c(:)) > + enddo > + !$acc end serial > + > + !$acc data copyin(readonly: a(:), b(:n)) copyin(c(:)) > + do i =3D 1,32 > + !$acc cache (readonly: a(:), b(:n)) > + !$acc cache (c(:)) > + enddo > + !$acc end data > + > + !$acc enter data copyin(readonly: a(:), b(:n)) copyin(c(:)) > + > +end subroutine foo > + > +program main > + integer :: g(32), h(32) > + integer :: i, n =3D 32, a(32) > + integer :: b(32), c(32) > + > + !$acc declare copyin(readonly: g), copyin(h) > + > + !$acc parallel copyin(readonly: a(:32), b(:n)) copyin(c(:)) > + do i =3D 1,32 > + !$acc cache (readonly: a(:), b(:n)) > + !$acc cache (c(:)) > + enddo > + !$acc end parallel > + > + !$acc kernels copyin(readonly: a(:), b(:n)) copyin(c(:)) > + do i =3D 1,32 > + !$acc cache (readonly: a(:), b(:n)) > + !$acc cache (c(:)) > + enddo > + !$acc end kernels > + > + !$acc serial copyin(readonly: a(:), b(:n)) copyin(c(:)) > + do i =3D 1,32 > + !$acc cache (readonly: a(:), b(:n)) > + !$acc cache (c(:)) > + enddo > + !$acc end serial > + > + !$acc data copyin(readonly: a(:), b(:n)) copyin(c(:)) > + do i =3D 1,32 > + !$acc cache (readonly: a(:), b(:n)) > + !$acc cache (c(:)) > + enddo > + !$acc end data > + > + !$acc enter data copyin(readonly: a(:), b(:n)) copyin(c(:)) > + > +end program main > + > +! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(read= only,to:\\*.+ map\\(alloc:a.+ map\\(readonly,to:\\*.+ map\\(alloc:b.+ map\\= (to:\\*.+ map\\(alloc:c.+" 1 "original" } } > +! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(read= only,to:a.+ map\\(alloc:a.+ map\\(readonly,to:b.+ map\\(alloc:b.+ map\\(to:= c.+ map\\(alloc:c.+" 1 "original" } } > +! { dg-final { scan-tree-dump-times "(?n)#pragma acc kernels map\\(reado= nly,to:\\*.+ map\\(alloc:a.+ map\\(readonly,to:\\*.+ map\\(alloc:b.+ map\\(= to:\\*.+ map\\(alloc:c.+" 1 "original" } } > +! { dg-final { scan-tree-dump-times "(?n)#pragma acc kernels map\\(reado= nly,to:a.+ map\\(alloc:a.+ map\\(readonly,to:b.+ map\\(alloc:b.+ map\\(to:c= .+ map\\(alloc:c.+" 1 "original" } } > +! { dg-final { scan-tree-dump-times "(?n)#pragma acc serial map\\(readon= ly,to:\\*.+ map\\(alloc:a.+ map\\(readonly,to:\\*.+ map\\(alloc:b.+ map\\(t= o:\\*.+ map\\(alloc:c.+" 1 "original" } } > +! { dg-final { scan-tree-dump-times "(?n)#pragma acc serial map\\(readon= ly,to:a.+ map\\(alloc:a.+ map\\(readonly,to:b.+ map\\(alloc:b.+ map\\(to:c.= + map\\(alloc:c.+" 1 "original" } } > +! { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(readonly= ,to:\\*.+ map\\(alloc:a.+ map\\(readonly,to:\\*.+ map\\(alloc:b.+ map\\(to:= \\*.+ map\\(alloc:c.+" 1 "original" } } > +! { dg-final { scan-tree-dump-times "(?n)#pragma acc data map\\(readonly= ,to:a.+ map\\(alloc:a.+ map\\(readonly,to:b.+ map\\(alloc:b.+ map\\(to:c.+ = map\\(alloc:c.+" 1 "original" } } > +! { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(re= adonly,to:\\*.+ map\\(alloc:a.+ map\\(readonly,to:\\*.+ map\\(alloc:b.+ map= \\(to:\\*.+ map\\(alloc:c.+" 1 "original" } } > +! { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(re= adonly,to:a.+ map\\(alloc:a.+ map\\(readonly,to:b.+ map\\(alloc:b.+ map\\(t= o:c.+ map\\(alloc:c.+" 1 "original" } } > + > +! { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(readonly:\= \*\\(integer\\(kind=3D4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\) \= \(readonly:\\*\\(integer\\(kind=3D4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len:= .+\\\]\\);" 8 "original" } } > +! { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(\\*\\(inte= ger\\(kind=3D4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\);" 8 "origi= nal" } } > diff --git a/gcc/tree-core.h b/gcc/tree-core.h > index 8a89462bd7e..d529712306d 100644 > --- a/gcc/tree-core.h > +++ b/gcc/tree-core.h > @@ -1344,6 +1344,12 @@ struct GTY(()) tree_base { > TYPE_READONLY in > all types >=20=20 > + OMP_CLAUSE_MAP_READONLY in > + OMP_CLAUSE_MAP > + > + OMP_CLAUSE__CACHE__READONLY in > + OMP_CLAUSE__CACHE_ > + > constant_flag: >=20=20 > TREE_CONSTANT in > diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc > index 654f5247e3a..926f7e006a7 100644 > --- a/gcc/tree-pretty-print.cc > +++ b/gcc/tree-pretty-print.cc > @@ -913,6 +913,8 @@ dump_omp_clause (pretty_printer *pp, tree clause, int= spc, dump_flags_t flags) >=20=20 > case OMP_CLAUSE_MAP: > pp_string (pp, "map("); > + if (OMP_CLAUSE_MAP_READONLY (clause)) > + pp_string (pp, "readonly,"); > switch (OMP_CLAUSE_MAP_KIND (clause)) > { > case GOMP_MAP_ALLOC: > @@ -1095,6 +1097,8 @@ dump_omp_clause (pretty_printer *pp, tree clause, i= nt spc, dump_flags_t flags) >=20=20 > case OMP_CLAUSE__CACHE_: > pp_string (pp, "("); > + if (OMP_CLAUSE__CACHE__READONLY (clause)) > + pp_string (pp, "readonly:"); > dump_generic_node (pp, OMP_CLAUSE_DECL (clause), > spc, flags, false); > goto print_clause_size; > diff --git a/gcc/tree.h b/gcc/tree.h > index e1fc6c2221d..b67a37d6522 100644 > --- a/gcc/tree.h > +++ b/gcc/tree.h > @@ -1841,6 +1841,14 @@ class auto_suppress_location_wrappers > #define OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE(NODE) \ > (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.addressable_fla= g) >=20=20 > +/* Nonzero if OpenACC 'readonly' modifier set, used for 'copyin'. */ > +#define OMP_CLAUSE_MAP_READONLY(NODE) \ > + TREE_READONLY (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)) > + > +/* Same as above, for use in OpenACC cache directives. */ > +#define OMP_CLAUSE__CACHE__READONLY(NODE) \ > + TREE_READONLY (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE__CACHE_)) > + > /* True on an OMP_CLAUSE_USE_DEVICE_PTR with an OpenACC 'if_present' > clause. */ > #define OMP_CLAUSE_USE_DEVICE_PTR_IF_PRESENT(NODE) \