public inbox for gcc-patches@gcc.gnu.org
 help / color / mirror / Atom feed
* [PATCH, OpenACC 2.7] readonly modifier support in front-ends
@ 2023-07-10 18:33 Chung-Lin Tang
  2023-07-11  7:00 ` Tobias Burnus
                   ` (2 more replies)
  0 siblings, 3 replies; 18+ messages in thread
From: Chung-Lin Tang @ 2023-07-10 18:33 UTC (permalink / raw)
  To: gcc-patches, Thomas Schwinge, Catherine Moore, Tobias Burnus

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

Hi Thomas,
this patch contains support for the 'readonly' modifier in copyin clauses
and the cache directive.

As we discussed earlier, the work for actually linking this to middle-end
points-to analysis is a somewhat non-trivial issue. This first patch allows
the language feature to be used in OpenACC directives first (with no effect for now).
The middle-end changes are probably going to be a later patch.

(Also CCing Tobias because of the Fortran bits)

Tested on powerpc64le-linux with nvptx offloading. Is this okay for trunk?

Thanks,
Chung-Lin

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

gcc/c/ChangeLog:
	* c-parser.cc (c_parser_omp_var_list_parens):
	Add 'bool *readonly = NULL' parameter, add readonly modifier parsing
	support.
	(c_parser_oacc_data_clause): Adjust c_parser_omp_var_list_parens call
	to turn on readonly modifier parsing for copyin clause, set
	OMP_CLAUSE_MAP_READONLY if readonly modifier found, update comments.
	(c_parser_oacc_cache): Adjust c_parser_omp_var_list_parens call
	to turn on readonly modifier parsing, set OMP_CLAUSE__CACHE__READONLY
	if readonly modifier found, update comments.

gcc/cp/ChangeLog:
	* parser.cc (cp_parser_omp_var_list):
	Add 'bool *readonly = NULL' parameter, add readonly modifier parsing
	support.
	(cp_parser_oacc_data_clause): Adjust cp_parser_omp_var_list call
	to turn on readonly modifier parsing for copyin clause, set
	OMP_CLAUSE_MAP_READONLY if readonly modifier found, update comments.
	(cp_parser_oacc_cache): Adjust cp_parser_omp_var_list call
	to turn on readonly modifier parsing, set OMP_CLAUSE__CACHE__READONLY
	if readonly modifier found, update comments.

gcc/fortran/ChangeLog:
	* gfortran.h (typedef struct gfc_omp_namelist): Adjust map_op as
	ENUM_BITFIELD field, add 'bool readonly' field.
	* openmp.cc (gfc_match_omp_map_clause): Add 'bool readonly = false'
	parameter, set n->u.readonly field.
	(gfc_match_omp_clauses): Add readonly modifier parsing for OpenACC
	copyin clause, adjust call to gfc_match_omp_map_clause.
	(gfc_match_oacc_cache): Add readonly modifier parsing for OpenACC
	cache directive, adjust call to gfc_match_omp_map_clause.
	* trans-openmp.cc (gfc_trans_omp_clauses): Set OMP_CLAUSE_MAP_READONLY,
	OMP_CLAUSE__CACHE__READONLY to 1 when readonly is set.

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: openacc-readonly-modifier.patch --]
[-- Type: text/plain, Size: 14732 bytes --]

diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index d4b98d5d8b6..09e1e89d793 100644
--- a/gcc/c/c-parser.cc
+++ 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)
 {
   /* The clauses location.  */
   location_t loc = c_parser_peek_token (parser)->location;
@@ -14067,6 +14068,20 @@ c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind,
   matching_parens parens;
   if (parens.require_open (parser))
     {
+      if (readonly != NULL)
+	{
+	  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;
+	    }
+	  else
+	    *readonly = false;
+	}
       list = c_parser_omp_variable_list (parser, loc, kind, list, allow_deref);
       parens.skip_until_found_close (parser);
     }
@@ -14084,7 +14099,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 +14154,22 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
     default:
       gcc_unreachable ();
     }
+
+  /* Turn on readonly modifier parsing for copyin clause.  */
+  bool readonly = false, *readonly_ptr = NULL;
+  if (c_kind == PRAGMA_OACC_CLAUSE_COPYIN)
+    readonly_ptr = &readonly;
+
   tree nl, c;
-  nl = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list, true);
+  nl = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list, true,
+				     readonly_ptr);
 
   for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
-    OMP_CLAUSE_SET_MAP_KIND (c, kind);
+    {
+      OMP_CLAUSE_SET_MAP_KIND (c, kind);
+      if (readonly)
+	OMP_CLAUSE_MAP_READONLY (c) = 1;
+    }
 
   return nl;
 }
@@ -18212,6 +18242,9 @@ 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.
 */
 
@@ -18219,8 +18252,14 @@ static tree
 c_parser_oacc_cache (location_t loc, c_parser *parser)
 {
   tree stmt, clauses;
+  bool readonly;
+
+  clauses = c_parser_omp_var_list_parens (parser, OMP_CLAUSE__CACHE_, NULL,
+					  false, &readonly);
+  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 acd1bd48af5..0f51289539b 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -37727,11 +37727,27 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
 
 static tree
 cp_parser_omp_var_list (cp_parser *parser, enum omp_clause_code kind, tree list,
-			bool allow_deref = false)
+			bool allow_deref = false, bool *readonly = NULL)
 {
   if (cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
-    return cp_parser_omp_var_list_no_open (parser, kind, list, NULL,
-					   allow_deref);
+    {
+      if (readonly != NULL)
+	{
+	  cp_token *token = cp_lexer_peek_token (parser->lexer);
+	  if (token->type == CPP_NAME
+	      && !strcmp (IDENTIFIER_POINTER (token->u.value), "readonly")
+	      && cp_lexer_nth_token_is (parser->lexer, 2, CPP_COLON))
+	    {
+	      cp_lexer_consume_token (parser->lexer);
+	      cp_lexer_consume_token (parser->lexer);
+	      *readonly = true;
+	    }
+	  else
+	    *readonly = false;
+	}
+      return cp_parser_omp_var_list_no_open (parser, kind, list, NULL,
+					     allow_deref);
+    }
   return list;
 }
 
@@ -37746,7 +37762,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,
@@ -37797,11 +37817,22 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
     default:
       gcc_unreachable ();
     }
+
+  /* Turn on readonly modifier parsing for copyin clause.  */
+  bool readonly = false, *readonly_ptr = NULL;
+  if (c_kind == PRAGMA_OACC_CLAUSE_COPYIN)
+    readonly_ptr = &readonly;
+
   tree nl, c;
-  nl = cp_parser_omp_var_list (parser, OMP_CLAUSE_MAP, list, true);
+  nl = cp_parser_omp_var_list (parser, OMP_CLAUSE_MAP, list, true,
+			       readonly_ptr);
 
   for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
-    OMP_CLAUSE_SET_MAP_KIND (c, kind);
+    {
+      OMP_CLAUSE_SET_MAP_KIND (c, kind);
+      if (readonly)
+	OMP_CLAUSE_MAP_READONLY (c) = 1;
+    }
 
   return nl;
 }
@@ -45875,6 +45906,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
@@ -45885,8 +45919,14 @@ cp_parser_oacc_cache (cp_parser *parser, cp_token *pragma_tok)
   auto_suppress_location_wrappers sentinel;
 
   tree stmt, clauses;
+  bool readonly;
+
+  clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE__CACHE_, NULL,
+				    false, &readonly);
+  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/gfortran.h b/gcc/fortran/gfortran.h
index cc7ba7c8846..9fa8962d63f 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) map_op:8;
+	  bool readonly;
+        };
       gfc_expr *align;
       struct
 	{
diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index 038907baa48..acd1428d2d7 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -1196,7 +1196,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,
@@ -1205,7 +1205,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.readonly = readonly;
+	}
       return true;
     }
 
@@ -2079,11 +2082,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],
@@ -4008,20 +4016,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.readonly = true;
+
   if (gfc_current_state() != COMP_DO 
       && gfc_current_state() != COMP_DO_CONCURRENT)
     {
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index 0f8323901d7..87d0b5e0cdf 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -3067,6 +3067,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 		      || (n->expr && gfc_expr_attr (n->expr).pointer)))
 		always_modifier = true;
 
+	      if (n->u.readonly)
+		OMP_CLAUSE_MAP_READONLY (node) = 1;
+
 	      switch (n->u.map_op)
 		{
 		case OMP_MAP_ALLOC:
@@ -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.readonly)
+		OMP_CLAUSE__CACHE__READONLY (node) = 1;
 	      omp_clauses = gfc_trans_add_clause (node, omp_clauses);
 	    }
 	  break;
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 a743e3cdfd8..6a9812c2253 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 3eebf5709b7..a79260e48eb 100644
--- a/gcc/tree.h
+++ 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_))
+
 /* True on an OMP_CLAUSE_USE_DEVICE_PTR with an OpenACC 'if_present'
    clause.  */
 #define OMP_CLAUSE_USE_DEVICE_PTR_IF_PRESENT(NODE) \

^ permalink raw reply	[flat|nested] 18+ messages in thread

end of thread, other threads:[~2024-05-16 12:36 UTC | newest]

Thread overview: 18+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2023-07-10 18:33 [PATCH, OpenACC 2.7] readonly modifier support in front-ends 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     ` [PATCH, OpenACC 2.7, v2] " Chung-Lin Tang
2023-10-26  9:43       ` 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

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).