public inbox for gcc-patches@gcc.gnu.org
 help / color / mirror / Atom feed
From: Chung-Lin Tang <chunglin.tang@siemens.com>
To: Tobias Burnus <tobias@codesourcery.com>,
	Thomas Schwinge <thomas@codesourcery.com>,
	Chung-Lin Tang <cltang@codesourcery.com>
Cc: gcc-patches@gcc.gnu.org, Catherine Moore <clm@codesourcery.com>
Subject: [PATCH, OpenACC 2.7, v2] readonly modifier support in front-ends
Date: Mon, 7 Aug 2023 21:58:27 +0800	[thread overview]
Message-ID: <b5af4407-1538-802f-92ca-aae843258c15@siemens.com> (raw)
In-Reply-To: <b45bd2bb-a306-5e42-7c84-53433e8d06a2@codesourcery.com>

[-- Attachment #1: Type: text/plain, Size: 6963 bytes --]

Hi Thomas, Tobias,
here's the updated v2 of the readonly modifier front-end patch.

On 2023/7/20 11:08 PM, Tobias Burnus wrote:
>>> +++ b/gcc/c/c-parser.cc
>>> @@ -14059,7 +14059,8 @@ c_parser_omp_variable_list (c_parser *parser,
>>>
>>>   static tree
>>>   c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind,
>>> -                           tree list, bool allow_deref = false)
>>> +                           tree list, bool allow_deref = false,
>>> +                           bool *readonly = NULL)
>>> ...
>> Instead of doing this in 'c_parser_omp_var_list_parens', I think it's
>> clearer to have this special 'readonly :' parsing logic in the two places
>> where it's used.
> I concur. The same issue also occurred for OpenMP's
> c_parser_omp_clause_to, and c_parser_omp_clause_from and the 'present'
> modifier. For it, I created a combined function but the main reason for
> that is that OpenMP also permits more modifiers (like 'iterators'),
> which would cause more duplication of code ('iterator' is not yet
> supported).
> 
> For something as simple to parse as this modifier, I would just do it at
> the two places – as Thomas suggested.

Okay, I've changed the C/C++ parser parts to have the parsing logic directly
added.

>>> +++ b/gcc/fortran/gfortran.h
>>> @@ -1360,7 +1360,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) map_op:8;
>>> +       bool readonly;
>>> +        };
>>>         gfc_expr *align;
>>>         struct
>>>        {
>> [...] Thus, the above looks good to me.
> I concur but I wonder whether it would be cleaner to name the struct;
> this makes it also more obvious what belongs together in the union.
> 
> Namely, naming the struct 'map' and then changing the 45 users from
> 'u.map_op' to 'u.map.op' and the new 'u.readonly' to 'u.map.readonly'. –
> this seems to be cleaner.

I've adjusted 'u.map' to be a named struct now, and updated the references.

>> + if (gfc_match ("readonly :") == MATCH_YES)
>> I note this one does not have a space after ':' in 'gfc_match', but the
>> one above in 'gfc_match_omp_clauses' does.  I don't know off-hand if that
>> makes a difference in parsing -- probably not, as all of
>> 'gcc/fortran/openmp.cc' generally doesn't seem to be very consistent
>> about these two variants?
> It *does* make a difference. And for obvious reasons. You don't want to permit:
> 
>    !$acc kernels asnyccopy(a)
> 
> but require at least one space (or comma) between "async" and "copy"..
> (In fixed form Fortran, it would be fine - as would be "!$acc k e nelsasy nc co p y(a)".)
> 
> A " " matches zero or more whitespaces, but with gfc_match_space you can find out
> whether there was whitespace or not.

Okay, made sure both are 'gfc_match ("readonly : ")'. Thanks for catching that, didn't
realize that space was significant.

>>> +++ 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.addressable_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

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 OMP_CLAUSE_MAP_READONLY == 1.

The other points-to patch then (also in front-ends) take the OMP_CLAUSE_MAP_READONLY
to mark the clauses of "base-pointers of array-sections" as OMP_CLAUSE_MAP_POINTS_TO_READONLY,
and later this gradually gets relayed to alias oracle routines in tree-ssa-alias.cc

Re-tested this v2 patch on powerpc64le-linux/nvptx. Okay for trunk?

Thanks,
Chung-Lin

2023-08-07  Chung-Lin Tang  <cltang@codesourcery.com>

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): Add 'bool readonly = false'
	parameter, set n->u.map.readonly field. Adjust 'n->u.map_op' to
	'n->u.map.op'.
	(gfc_match_omp_clause_reduction): Adjust 'n->u.map_op' to 'n->u.map.op'.
	(gfc_match_omp_clauses): Add readonly modifier parsing for OpenACC
	copyin clause, adjust call to gfc_match_omp_map_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-pretty-print.cc (dump_omp_clause): Add support for printing
	OMP_CLAUSE_MAP_READONLY and OMP_CLAUSE__CACHE__READONLY.
	* tree.h (OMP_CLAUSE_MAP_READONLY): New macro.
	(OMP_CLAUSE__CACHE__READONLY): New macro.

gcc/testsuite/ChangeLog:
	* c-c++-common/goacc/readonly-1.c: New test.
	* gfortran.dg/goacc/readonly-1.f90: New test.




[-- Attachment #2: readonly-fe-v2.patch --]
[-- Type: text/plain, Size: 26539 bytes --]

diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 24a6eb6e459..5779f499ae1 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -14084,7 +14084,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 )
+ */
 
 static tree
 c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
@@ -14135,11 +14139,36 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
     default:
       gcc_unreachable ();
     }
-  tree nl, c;
-  nl = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list, true);
 
-  for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
-    OMP_CLAUSE_SET_MAP_KIND (c, kind);
+  tree nl = list;
+  bool readonly = false;
+  matching_parens parens;
+  if (parens.require_open (parser))
+    {
+      /* Turn on readonly modifier parsing for copyin clause.  */
+      if (c_kind == PRAGMA_OACC_CLAUSE_COPYIN)
+	{
+	  c_token *token = c_parser_peek_token (parser);
+	  if (token->type == CPP_NAME
+	      && !strcmp (IDENTIFIER_POINTER (token->value), "readonly")
+	      && c_parser_peek_2nd_token (parser)->type == CPP_COLON)
+	    {
+	      c_parser_consume_token (parser);
+	      c_parser_consume_token (parser);
+	      readonly = true;
+	    }
+	}
+      location_t loc = c_parser_peek_token (parser)->location;
+      nl = c_parser_omp_variable_list (parser, loc, OMP_CLAUSE_MAP, list, true);
+      parens.skip_until_found_close (parser);
+    }
+
+  for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
+    {
+      OMP_CLAUSE_SET_MAP_KIND (c, kind);
+      if (readonly)
+	OMP_CLAUSE_MAP_READONLY (c) = 1;
+    }
 
   return nl;
 }
@@ -18161,15 +18190,40 @@ c_parser_omp_structured_block (c_parser *parser, bool *if_p)
 /* OpenACC 2.0:
    # pragma acc cache (variable-list) new-line
 
+   OpenACC 2.7:
+   # pragma acc cache (readonly: variable-list) new-line
+
    LOC is the location of the #pragma token.
 */
 
 static tree
 c_parser_oacc_cache (location_t loc, c_parser *parser)
 {
-  tree stmt, clauses;
+  tree stmt, clauses = NULL_TREE;
+  bool readonly = false;
+  matching_parens parens;
+
+  if (parens.require_open (parser))
+    {
+      c_token *token = c_parser_peek_token (parser);
+      if (token->type == CPP_NAME
+	  && !strcmp (IDENTIFIER_POINTER (token->value), "readonly")
+	  && c_parser_peek_2nd_token (parser)->type == CPP_COLON)
+	{
+	  c_parser_consume_token (parser);
+	  c_parser_consume_token (parser);
+	  readonly = true;
+	}
+      location_t loc = c_parser_peek_token (parser)->location;
+      clauses = c_parser_omp_variable_list (parser, loc, OMP_CLAUSE__CACHE_,
+					    NULL_TREE);
+      parens.skip_until_found_close (parser);
+    }
+
+  if (readonly)
+    for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+      OMP_CLAUSE__CACHE__READONLY (c) = 1;
 
-  clauses = c_parser_omp_var_list_parens (parser, OMP_CLAUSE__CACHE_, NULL);
   clauses = c_finish_omp_clauses (clauses, C_ORT_ACC);
 
   c_parser_skip_to_pragma_eol (parser);
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index d7ef5b34d42..ac8a656874a 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -37750,7 +37750,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 )
+ */
 
 static tree
 cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
@@ -37801,11 +37805,33 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
     default:
       gcc_unreachable ();
     }
-  tree nl, c;
-  nl = cp_parser_omp_var_list (parser, OMP_CLAUSE_MAP, list, true);
+  tree nl = list;
+  bool readonly = false;
+  if (cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
+    {
+      /* Turn on readonly modifier parsing for copyin clause.  */
+      if (c_kind == PRAGMA_OACC_CLAUSE_COPYIN)
+	{
+	  cp_token *token = cp_lexer_peek_token (parser->lexer);
+	  if (token->type == CPP_NAME
+	      && !strcmp (IDENTIFIER_POINTER (token->u.value), "readonly")
+	      && cp_lexer_peek_nth_token (parser->lexer, 2)->type == CPP_COLON)
+	    {
+	      cp_lexer_consume_token (parser->lexer);
+	      cp_lexer_consume_token (parser->lexer);
+	      readonly = true;
+	    }
+	}
+      nl = cp_parser_omp_var_list_no_open (parser, OMP_CLAUSE_MAP, list, NULL,
+					   true);
+    }
 
-  for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
-    OMP_CLAUSE_SET_MAP_KIND (c, kind);
+  for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
+    {
+      OMP_CLAUSE_SET_MAP_KIND (c, kind);
+      if (readonly)
+	OMP_CLAUSE_MAP_READONLY (c) = 1;
+    }
 
   return nl;
 }
@@ -45825,6 +45851,9 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
 
 /* OpenACC 2.0:
    # pragma acc cache (variable-list) new-line
+
+   OpenACC 2.7:
+   # pragma acc cache (readonly: variable-list) new-line
 */
 
 static tree
@@ -45834,9 +45863,28 @@ cp_parser_oacc_cache (cp_parser *parser, cp_token *pragma_tok)
      clauses.  */
   auto_suppress_location_wrappers sentinel;
 
-  tree stmt, clauses;
+  tree stmt, clauses = NULL_TREE;
+  bool readonly = false;
+
+  if (cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
+    {
+      cp_token *token = cp_lexer_peek_token (parser->lexer);
+      if (token->type == CPP_NAME
+	  && !strcmp (IDENTIFIER_POINTER (token->u.value), "readonly")
+	  && cp_lexer_peek_nth_token (parser->lexer, 2)->type == CPP_COLON)
+	{
+	  cp_lexer_consume_token (parser->lexer);
+	  cp_lexer_consume_token (parser->lexer);
+	  readonly = true;
+	}
+      clauses = cp_parser_omp_var_list_no_open (parser, OMP_CLAUSE__CACHE_,
+						NULL, NULL);
+    }
+
+  if (readonly)
+    for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+      OMP_CLAUSE__CACHE__READONLY (c) = 1;
 
-  clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE__CACHE_, NULL_TREE);
   clauses = finish_omp_clauses (clauses, C_ORT_ACC);
 
   cp_parser_require_pragma_eol (parser, cp_lexer_peek_token (parser->lexer));
diff --git a/gcc/fortran/dump-parse-tree.cc b/gcc/fortran/dump-parse-tree.cc
index 68122e3e6fd..0e888fafe7b 100644
--- a/gcc/fortran/dump-parse-tree.cc
+++ b/gcc/fortran/dump-parse-tree.cc
@@ -1398,6 +1398,9 @@ show_omp_namelist (int list_type, gfc_omp_namelist *n)
 	    fputs (") ALLOCATE(", dumpfile);
 	  continue;
 	}
+      if ((list_type == OMP_LIST_MAP || list_type == OMP_LIST_CACHE)
+	  && n->u.map.readonly)
+	fputs ("readonly,", dumpfile);
       if (list_type == OMP_LIST_REDUCTION)
 	switch (n->u.reduction_op)
 	  {
@@ -1465,7 +1468,7 @@ show_omp_namelist (int list_type, gfc_omp_namelist *n)
 	  default: break;
 	  }
       else if (list_type == 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 9a00e6dea6f..a8667a6e6d3 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -1360,7 +1360,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 2952cd300ac..af769d9efbd 100644
--- 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 readonly = false)
 {
   gfc_omp_namelist **head = 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 **list, gfc_omp_map_op map_op,
     {
       gfc_omp_namelist *n;
       for (n = *head; n; n = n->next)
-	n->u.map_op = map_op;
+	{
+	  n->u.map.op = map_op;
+	  n->u.map.readonly = readonly;
+	}
       return true;
     }
 
@@ -1520,7 +1523,7 @@ gfc_match_omp_clause_reduction (char pc, gfc_omp_clauses *c, bool openacc,
 	    gfc_omp_namelist *p = gfc_get_omp_namelist (), **tl;
 	    p->sym = n->sym;
 	    p->where = p->where;
-	    p->u.map_op = OMP_MAP_ALWAYS_TOFROM;
+	    p->u.map.op = OMP_MAP_ALWAYS_TOFROM;
 
 	    tl = &c->lists[OMP_LIST_MAP];
 	    while (*tl)
@@ -2180,11 +2183,16 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
 	    {
 	      if (openacc)
 		{
-		  if (gfc_match ("copyin ( ") == MATCH_YES
-		      && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
-						   OMP_MAP_TO, true,
-						   allow_derived))
-		    continue;
+		  if (gfc_match ("copyin ( ") == MATCH_YES)
+		    {
+		      bool readonly = false;
+		      if (gfc_match ("readonly : ") == MATCH_YES)
+			readonly = true;
+		      if (gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
+						    OMP_MAP_TO, true,
+						    allow_derived, readonly))
+			continue;
+		    }
 		}
 	      else if (gfc_match_omp_variable_list ("copyin (",
 						    &c->lists[OMP_LIST_COPYIN],
@@ -3101,7 +3109,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
 		{
 		  gfc_omp_namelist *n;
 		  for (n = *head; n; n = n->next)
-		    n->u.map_op = map_op;
+		    n->u.map.op = map_op;
 		  continue;
 		}
 	      gfc_current_locus = old_loc;
@@ -3942,7 +3950,7 @@ gfc_match_oacc_declare (void)
       if (gfc_current_ns->proc_name
 	  && gfc_current_ns->proc_name->attr.flavor == FL_MODULE)
 	{
-	  if (n->u.map_op != OMP_MAP_ALLOC && n->u.map_op != OMP_MAP_TO)
+	  if (n->u.map.op != OMP_MAP_ALLOC && n->u.map.op != OMP_MAP_TO)
 	    {
 	      gfc_error ("Invalid clause in module with !$ACC DECLARE at %L",
 			 &where);
@@ -3976,7 +3984,7 @@ gfc_match_oacc_declare (void)
 	  return MATCH_ERROR;
 	}
 
-      switch (n->u.map_op)
+      switch (n->u.map.op)
 	{
 	  case OMP_MAP_FORCE_ALLOC:
 	  case OMP_MAP_ALLOC:
@@ -4091,20 +4099,35 @@ gfc_match_oacc_wait (void)
 match
 gfc_match_oacc_cache (void)
 {
+  bool readonly = false;
   gfc_omp_clauses *c = 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 like a
      only_sections variant next to its allow_sections parameter.  */
-  match m = gfc_match_omp_variable_list (" (",
-					 &c->lists[OMP_LIST_CACHE], true,
-					 NULL, NULL, true);
+  match m = gfc_match (" ( ");
   if (m != MATCH_YES)
     {
       gfc_free_omp_clauses(c);
       return m;
     }
 
+  if (gfc_match ("readonly : ") == MATCH_YES)
+    readonly = true;
+
+  gfc_omp_namelist **head = NULL;
+  m = gfc_match_omp_variable_list ("", &c->lists[OMP_LIST_CACHE], true,
+				   NULL, &head, true);
+  if (m != MATCH_YES)
+    {
+      gfc_free_omp_clauses(c);
+      return m;
+    }
+
+  if (readonly)
+    for (gfc_omp_namelist *n = *head; n; n = n->next)
+      n->u.map.readonly = true;
+
   if (gfc_current_state() != COMP_DO 
       && gfc_current_state() != COMP_DO_CONCURRENT)
     {
@@ -8142,8 +8165,8 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 		  }
 		if (openacc
 		    && list == OMP_LIST_MAP
-		    && (n->u.map_op == OMP_MAP_ATTACH
-			|| n->u.map_op == OMP_MAP_DETACH))
+		    && (n->u.map.op == OMP_MAP_ATTACH
+			|| n->u.map.op == OMP_MAP_DETACH))
 		  {
 		    symbol_attribute attr;
 		    if (n->expr)
@@ -8153,7 +8176,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 		    if (!attr.pointer && !attr.allocatable)
 		      gfc_error ("%qs clause argument must be ALLOCATABLE or "
 				 "a POINTER at %L",
-				 (n->u.map_op == OMP_MAP_ATTACH) ? "attach"
+				 (n->u.map.op == OMP_MAP_ATTACH) ? "attach"
 				 : "detach", &n->where);
 		  }
 		if (lastref
@@ -8224,7 +8247,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 		else if (openacc)
 		  {
 		    if (list == OMP_LIST_MAP
-			&& n->u.map_op == OMP_MAP_FORCE_DEVICEPTR)
+			&& n->u.map.op == OMP_MAP_FORCE_DEVICEPTR)
 		      resolve_oacc_deviceptr_clause (n->sym, n->where, name);
 		    else
 		      resolve_oacc_data_clauses (n->sym, n->where, name);
@@ -8246,7 +8269,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *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:
@@ -8273,7 +8296,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *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:
@@ -8283,16 +8306,16 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 			case OMP_MAP_PRESENT_ALLOC:
 			  break;
 			case OMP_MAP_TOFROM:
-			  n->u.map_op = OMP_MAP_TO;
+			  n->u.map.op = OMP_MAP_TO;
 			  break;
 			case OMP_MAP_ALWAYS_TOFROM:
-			  n->u.map_op = OMP_MAP_ALWAYS_TO;
+			  n->u.map.op = OMP_MAP_ALWAYS_TO;
 			  break;
 			case OMP_MAP_PRESENT_TOFROM:
-			  n->u.map_op = OMP_MAP_PRESENT_TO;
+			  n->u.map.op = OMP_MAP_PRESENT_TO;
 			  break;
 			case OMP_MAP_ALWAYS_PRESENT_TOFROM:
-			  n->u.map_op = OMP_MAP_ALWAYS_PRESENT_TO;
+			  n->u.map.op = OMP_MAP_ALWAYS_PRESENT_TO;
 			  break;
 			default:
 			  gfc_error ("TARGET ENTER DATA with map-type other "
@@ -8302,7 +8325,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *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:
@@ -8312,16 +8335,16 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 			case OMP_MAP_DELETE:
 			  break;
 			case OMP_MAP_TOFROM:
-			  n->u.map_op = OMP_MAP_FROM;
+			  n->u.map.op = OMP_MAP_FROM;
 			  break;
 			case OMP_MAP_ALWAYS_TOFROM:
-			  n->u.map_op = OMP_MAP_ALWAYS_FROM;
+			  n->u.map.op = OMP_MAP_ALWAYS_FROM;
 			  break;
 			case OMP_MAP_PRESENT_TOFROM:
-			  n->u.map_op = OMP_MAP_PRESENT_FROM;
+			  n->u.map.op = OMP_MAP_PRESENT_FROM;
 			  break;
 			case OMP_MAP_ALWAYS_PRESENT_TOFROM:
-			  n->u.map_op = OMP_MAP_ALWAYS_PRESENT_FROM;
+			  n->u.map.op = 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 b0fd25e92a3..1ff1dda026a 100644
--- a/gcc/fortran/trans-decl.cc
+++ b/gcc/fortran/trans-decl.cc
@@ -6614,7 +6614,7 @@ add_clause (gfc_symbol *sym, gfc_omp_map_op map_op)
 
   n = gfc_get_omp_namelist ();
   n->sym = sym;
-  n->u.map_op = map_op;
+  n->u.map.op = map_op;
 
   if (!module_oacc_clauses)
     module_oacc_clauses = gfc_get_omp_clauses ();
@@ -6716,10 +6716,10 @@ finish_oacc_declare (gfc_namespace *ns, gfc_symbol *sym, bool block)
 
   for (n = omp_clauses->lists[OMP_LIST_MAP]; n; n = n->next)
     {
-      switch (n->u.map_op)
+      switch (n->u.map.op)
 	{
 	  case OMP_MAP_DEVICE_RESIDENT:
-	    n->u.map_op = OMP_MAP_FORCE_ALLOC;
+	    n->u.map.op = OMP_MAP_FORCE_ALLOC;
 	    break;
 
 	  default:
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index cf741cebf91..a4628e460bd 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -3067,7 +3067,10 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		      || (n->expr && gfc_expr_attr (n->expr).pointer)))
 		always_modifier = true;
 
-	      switch (n->u.map_op)
+	      if (n->u.map.readonly)
+		OMP_CLAUSE_MAP_READONLY (node) = 1;
+
+	      switch (n->u.map.op)
 		{
 		case OMP_MAP_ALLOC:
 		  OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALLOC);
@@ -3194,8 +3197,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 != EXEC_OMP_TARGET_EXIT_DATA
-		      && n->u.map_op != OMP_MAP_DELETE
-		      && n->u.map_op != OMP_MAP_RELEASE)
+		      && n->u.map.op != OMP_MAP_DELETE
+		      && n->u.map.op != OMP_MAP_RELEASE)
 		    {
 		      gcc_assert (n->sym->ts.u.cl->backend_decl);
 		      node5 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
@@ -3261,7 +3264,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 			{
 			  enum gomp_map_kind gmk = GOMP_MAP_POINTER;
 			  if (op == EXEC_OMP_TARGET_EXIT_DATA
-			      && n->u.map_op == OMP_MAP_DELETE)
+			      && n->u.map.op == OMP_MAP_DELETE)
 			    gmk = GOMP_MAP_DELETE;
 			  else if (op == EXEC_OMP_TARGET_EXIT_DATA)
 			    gmk = GOMP_MAP_RELEASE;
@@ -3284,7 +3287,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 			{
 			  enum gomp_map_kind gmk;
 			  if (op == EXEC_OMP_TARGET_EXIT_DATA
-			      && n->u.map_op == OMP_MAP_DELETE)
+			      && n->u.map.op == OMP_MAP_DELETE)
 			    gmk = GOMP_MAP_DELETE;
 			  else if (op == EXEC_OMP_TARGET_EXIT_DATA)
 			    gmk = GOMP_MAP_RELEASE;
@@ -3316,18 +3319,18 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		      node2 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
 		      OMP_CLAUSE_DECL (node2) = decl;
 		      OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
-		      if (n->u.map_op == OMP_MAP_DELETE)
+		      if (n->u.map.op == OMP_MAP_DELETE)
 			map_kind = GOMP_MAP_DELETE;
 		      else if (op == EXEC_OMP_TARGET_EXIT_DATA
-			       || n->u.map_op == OMP_MAP_RELEASE)
+			       || n->u.map.op == OMP_MAP_RELEASE)
 			map_kind = GOMP_MAP_RELEASE;
 		      else
 			map_kind = GOMP_MAP_TO_PSET;
 		      OMP_CLAUSE_SET_MAP_KIND (node2, map_kind);
 
 		      if (op != EXEC_OMP_TARGET_EXIT_DATA
-			  && n->u.map_op != OMP_MAP_DELETE
-			  && n->u.map_op != OMP_MAP_RELEASE)
+			  && n->u.map.op != OMP_MAP_DELETE
+			  && n->u.map.op != OMP_MAP_RELEASE)
 			{
 			  node3 = build_omp_clause (input_location,
 						    OMP_CLAUSE_MAP);
@@ -3345,7 +3348,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 			      = gfc_conv_descriptor_data_get (decl);
 			  OMP_CLAUSE_SIZE (node3) = size_int (0);
 
-			  if (n->u.map_op == OMP_MAP_ATTACH)
+			  if (n->u.map.op == OMP_MAP_ATTACH)
 			    {
 			      /* Standalone attach clauses used with arrays with
 				 descriptors must copy the descriptor to the
@@ -3361,7 +3364,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 			      node3 = NULL;
 			      goto finalize_map_clause;
 			    }
-			  else if (n->u.map_op == OMP_MAP_DETACH)
+			  else if (n->u.map.op == OMP_MAP_DETACH)
 			    {
 			      OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_DETACH);
 			      /* Similarly to above, we don't want to unmap PTR
@@ -3553,8 +3556,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 == OMP_MAP_ATTACH
-			      || n->u.map_op == OMP_MAP_DETACH))
+			  && (n->u.map.op == OMP_MAP_ATTACH
+			      || n->u.map.op == OMP_MAP_DETACH))
 			{
 			  OMP_CLAUSE_DECL (node)
 			    = build_fold_addr_expr (OMP_CLAUSE_DECL (node));
@@ -3585,7 +3588,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 					   fold_convert (size_type_node,
 					       se.string_length),
 					   TYPE_SIZE_UNIT (tmp));
-			  if (n->u.map_op == OMP_MAP_DELETE)
+			  if (n->u.map.op == OMP_MAP_DELETE)
 			    kind = GOMP_MAP_DELETE;
 			  else if (op == EXEC_OMP_TARGET_EXIT_DATA)
 			    kind = GOMP_MAP_RELEASE;
@@ -3642,8 +3645,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 == OMP_MAP_ATTACH
-				  || n->u.map_op == OMP_MAP_DETACH))
+			      && (n->u.map.op == OMP_MAP_ATTACH
+				  || n->u.map.op == OMP_MAP_DETACH))
 			    {
 			      OMP_CLAUSE_DECL (node)
 				= build_fold_addr_expr (inner);
@@ -3689,8 +3692,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 == OMP_MAP_ATTACH
-			   || n->u.map_op == OMP_MAP_DETACH)
+		      if ((n->u.map.op == OMP_MAP_ATTACH
+			   || n->u.map.op == OMP_MAP_DETACH)
 			  && (POINTER_TYPE_P (TREE_TYPE (inner))
 			      || GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (inner))))
 			{
@@ -3724,8 +3727,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 			    map_kind = ((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 == OMP_MAP_RELEASE
-				   || n->u.map_op == OMP_MAP_DELETE)
+			  else if (n->u.map.op == OMP_MAP_RELEASE
+				   || n->u.map.op == OMP_MAP_DELETE)
 			    ;
 			  else if (op == EXEC_OMP_TARGET_EXIT_DATA)
 			    map_kind = GOMP_MAP_RELEASE;
@@ -3920,6 +3923,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		}
 	      if (n->u.present_modifier)
 		OMP_CLAUSE_MOTION_PRESENT (node) = 1;
+	      if (list == OMP_LIST_CACHE && n->u.map.readonly)
+		OMP_CLAUSE__CACHE__READONLY (node) = 1;
 	      omp_clauses = gfc_trans_add_clause (node, omp_clauses);
 	    }
 	  break;
@@ -6333,7 +6338,7 @@ gfc_add_clause_implicitly (gfc_omp_clauses *clauses_out,
 	  n2->where = n->where;
 	  n2->sym = n->sym;
 	  if (is_target)
-	    n2->u.map_op = OMP_MAP_TOFROM;
+	    n2->u.map.op = OMP_MAP_TOFROM;
 	  if (tail)
 	    {
 	      tail->next = n2;
diff --git a/gcc/testsuite/c-c++-common/goacc/readonly-1.c b/gcc/testsuite/c-c++-common/goacc/readonly-1.c
new file mode 100644
index 00000000000..171f96c08db
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/readonly-1.c
@@ -0,0 +1,27 @@
+/* { dg-additional-options "-fdump-tree-original" } */
+
+struct S
+{
+  int *ptr;
+  float f;
+};
+
+
+int main (void)
+{
+  int x[32];
+  struct S s = {x, 0};
+
+  #pragma acc parallel copyin(readonly: x[:32], s.ptr[:16])
+  {
+    #pragma acc cache (readonly: x[:32])
+  }
+  return 0;
+}
+
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:\\*s.ptr \\\[len: 64\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: 128\\\]\\)" 1 "original" { target { c } } } } */
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:\\*NON_LVALUE_EXPR <s.ptr> \\\[len: 64\\\]\\) .+ map\\(readonly,to:x\\\[0\\\] \\\[len: 128\\\]\\)" 1 "original" { target { c++ } } } } */
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(readonly:x\\\[0\\\] \\\[len: 128\\\]\\);$" 1 "original" } } */
+
+
+
diff --git a/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90 b/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90
new file mode 100644
index 00000000000..069fec0a0d5
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/readonly-1.f90
@@ -0,0 +1,28 @@
+! { dg-additional-options "-fdump-tree-original" }
+
+subroutine foo (a, n)
+  integer :: n, a(:)
+  integer :: i, b(n)
+  !$acc parallel copyin(readonly: a(:), b(:n))
+  do i = 1,32
+     !$acc cache (readonly: a(:), b(:n))
+  enddo
+  !$acc end parallel
+end subroutine foo
+
+program main
+  integer :: i, n = 32, a(32)
+  integer :: b(32)
+  !$acc parallel copyin(readonly: a(:32), b(:n))
+  do i = 1,32
+     !$acc cache (readonly: a(:), b(:n))
+  enddo
+  !$acc end parallel
+end program main
+
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\) .+ map\\(readonly,to:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc parallel map\\(readonly,to:a\\\[\\(\\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) &a\\) / 4\\\] \\\[len: .+\\\]\\) .+ map\\(readonly,to:b\\\[\\(\\(integer\\(kind=8\\)\\) parm.*data - \\(integer\\(kind=8\\)\\) &b\\) / 4\\\] \\\[len: .+\\\]\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(readonly:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\) \\(readonly:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\);" 2 "original" } }
+
+
+
diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc
index 25d191b10fd..9604c3eecc5 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -905,6 +905,8 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
 
     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:
@@ -1075,6 +1077,8 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
 
     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 4c04245e2b1..1301491587f 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1811,6 +1811,14 @@ class auto_suppress_location_wrappers
 #define OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE(NODE) \
   (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.addressable_flag)
 
+/* Nonzero if OpenACC 'readonly' modifier set, used for 'copyin'.  */
+#define OMP_CLAUSE_MAP_READONLY(NODE) \
+  TREE_NOTHROW (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP))
+
+/* Same as above, for use in OpenACC cache directives.  */
+#define OMP_CLAUSE__CACHE__READONLY(NODE) \
+  TREE_NOTHROW (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) \

  reply	other threads:[~2023-08-07 13:58 UTC|newest]

Thread overview: 18+ messages / expand[flat|nested]  mbox.gz  Atom feed  top
2023-07-10 18:33 [PATCH, OpenACC 2.7] " Chung-Lin Tang
2023-07-11  7:00 ` Tobias Burnus
2023-07-20 13:33 ` Thomas Schwinge
2023-07-20 15:08   ` Tobias Burnus
2023-08-07 13:58     ` Chung-Lin Tang [this message]
2023-10-26  9:43       ` [PATCH, OpenACC 2.7, v2] " Thomas Schwinge
2024-03-07  8:02         ` Chung-Lin Tang
2024-03-13  9:12           ` Thomas Schwinge
2024-03-14 15:09             ` OpenACC 2.7: front-end support for readonly modifier: Add basic OpenACC 'declare' testing (was: [PATCH, OpenACC 2.7, v2] readonly modifier support in front-ends) Thomas Schwinge
2024-03-14 16:55               ` OpenACC 2.7: front-end support for readonly modifier: Add basic OpenACC 'declare' testing Tobias Burnus
2024-03-14 16:55               ` Tobias Burnus
2023-07-25 15:52 ` [PATCH, OpenACC 2.7] Connect readonly modifier to points-to analysis Chung-Lin Tang
2023-10-27 14:28   ` Thomas Schwinge
2023-10-30 12:46     ` Richard Biener
2024-04-03 11:50       ` Chung-Lin Tang
2024-04-11 14:29         ` Thomas Schwinge
2024-04-12  6:17           ` Richard Biener
2024-05-16 12:36         ` Richard Biener

Reply instructions:

You may reply publicly to this message via plain-text email
using any one of the following methods:

* Save the following mbox file, import it into your mail client,
  and reply-to-all from there: mbox

  Avoid top-posting and favor interleaved quoting:
  https://en.wikipedia.org/wiki/Posting_style#Interleaved_style

* Reply using the --to, --cc, and --in-reply-to
  switches of git-send-email(1):

  git send-email \
    --in-reply-to=b5af4407-1538-802f-92ca-aae843258c15@siemens.com \
    --to=chunglin.tang@siemens.com \
    --cc=clm@codesourcery.com \
    --cc=cltang@codesourcery.com \
    --cc=gcc-patches@gcc.gnu.org \
    --cc=thomas@codesourcery.com \
    --cc=tobias@codesourcery.com \
    /path/to/YOUR_REPLY

  https://kernel.org/pub/software/scm/git/docs/git-send-email.html

* If your mail client supports setting the In-Reply-To header
  via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line before the message body.
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).