OpenACC 2.7: front-end support for readonly modifier

This patch implements the front-end support for the 'readonly' modifier for the
OpenACC 'copyin' clause and 'cache' directive.

This currently only includes front-end parsing for C/C++/Fortran and setting of
new bits OMP_CLAUSE_MAP_READONLY, OMP_CLAUSE__CACHE__READONLY. Further linking
of these bits to points-to analysis and/or utilization of read-only memory in
accelerator target are for later patches.

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 53e99aa29d9695d4e32abb5c8f6527fce1c1b8e2..00f8bf4376e537e04ea8e468a05dade3c7212d8b 100644 (file)
--- 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 )
+ */
 
 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, pragma_omp_clause c_kind,
     default:
       gcc_unreachable ();
     }
-  tree nl, c;
-  nl = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list, false);
 
-  for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
-    OMP_CLAUSE_SET_MAP_KIND (c, kind);
+  tree nl = list;
+  bool readonly = false;
+  location_t open_loc = 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 == 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;
+           }
+       }
+      nl = c_parser_omp_variable_list (parser, open_loc, OMP_CLAUSE_MAP, list,
+                                      false);
+      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;
 }
@@ -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
 
+   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;
+  location_t open_loc = c_parser_peek_token (parser)->location;
+  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;
+       }
+      clauses = c_parser_omp_variable_list (parser, open_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 bc3aa9dd6adb562721211e92bd9f3aa4e5d9a06d..6134187c5db92681c64078211fc0f309d8ac87c0 100644 (file)
--- 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 )
+ */
 
 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 = cp_parser_omp_var_list (parser, OMP_CLAUSE_MAP, list, false);
 
-  for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
-    OMP_CLAUSE_SET_MAP_KIND (c, kind);
+  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,
+                                          false);
+    }
+
+  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;
 }
@@ -47178,6 +47205,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
@@ -47187,9 +47217,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 7b154eb3ca7193927913a96358a0d211ab5b6240..db84b06289be047743ef55186fd51afd07d18e05 100644 (file)
--- 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 == 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)
          {
@@ -1467,7 +1470,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 ebba2336e129378ea537a7a209a9d5721c3b7979..32b792f85fb900540e366cddb2d5d33d0c4fed90 100644 (file)
--- 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 38de60238c0e01bfb6a2fa083352854a2e8db83d..5c44e666eb998181cfd4fa349199dc5ff6e156de 100644 (file)
--- 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 = *head; n; n = n->next)
-       n->u.map_op = map_op;
+       n->u.map.op = map_op;
       return true;
     }
 
@@ -1524,7 +1524,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)
@@ -2181,11 +2181,25 @@ 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 = gfc_match ("readonly : ") == MATCH_YES;
+                     head = NULL;
+                     if (gfc_match_omp_variable_list ("",
+                                                      &c->lists[OMP_LIST_MAP],
+                                                      true, NULL, &head, true,
+                                                      allow_derived)
+                         == MATCH_YES)
+                       {
+                         gfc_omp_namelist *n;
+                         for (n = *head; n; n = n->next)
+                           {
+                             n->u.map.op = OMP_MAP_TO;
+                             n->u.map.readonly = 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 = *head; n; n = n->next)
-                   n->u.map_op = map_op;
+                   n->u.map.op = map_op;
                  continue;
                }
              gfc_current_locus = old_loc;
@@ -4002,7 +4016,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);
@@ -4036,7 +4050,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:
@@ -4151,21 +4165,36 @@ 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_current_state() != COMP_DO 
+  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)
     {
       gfc_error ("ACC CACHE directive must be inside of loop %C");
@@ -8436,8 +8465,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)
@@ -8447,7 +8476,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
@@ -8518,7 +8547,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);
@@ -8540,7 +8569,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:
@@ -8567,7 +8596,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:
@@ -8577,16 +8606,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 "
@@ -8596,7 +8625,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:
@@ -8606,16 +8635,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 6d4630369660b79b728a8646da086f600e1492ad..b7dea11461fb6c214863200dca0f952b0ecf3155 100644 (file)
--- 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)
 
   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 ();
@@ -6846,10 +6846,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 1dba47126edce7fb2de5ddddf88c2933813a8a0e..7a088ec5b2dfd7756d9b26fc97d9852c0da51d7d 100644 (file)
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -3148,7 +3148,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);
@@ -3275,8 +3278,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);
@@ -3342,7 +3345,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;
@@ -3365,7 +3368,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;
@@ -3397,18 +3400,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);
@@ -3426,7 +3429,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
@@ -3442,7 +3445,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
@@ -3635,8 +3638,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));
@@ -3665,7 +3668,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 == 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;
@@ -3722,8 +3725,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);
@@ -3815,8 +3818,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))))
                        {
@@ -3849,8 +3852,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
                                   || op == EXEC_OACC_EXIT_DATA)
@@ -4097,6 +4100,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;
@@ -6570,7 +6575,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 (file)
index 0000000..34fc92c
--- /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 = {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 "original" { target { c } } } } */
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc enter 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 "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 <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:\\*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 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 <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 enter data 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 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/testsuite/gfortran.dg/goacc/readonly-1.f90
new file mode 100644 (file)
index 0000000..696ebd0
--- /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 = 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 = 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 = 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 = 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 = 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 = 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 = 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 = 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 = 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\\(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 parallel 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 kernels 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 kernels 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 serial 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 serial 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 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\\(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 enter 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 cache \\(readonly:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\) \\(readonly:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\);" 8 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc cache \\(\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) parm.*data \\\[len: .+\\\]\\);" 8 "original" } }

diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index 8a89462bd7ecac52fcdc11c0b57ccf7c190572b3..d529712306d9a0592d8326077e20da70a58fefe6 100644 (file)
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -1344,6 +1344,12 @@ struct GTY(()) tree_base {
        TYPE_READONLY in
            all types
 
+       OMP_CLAUSE_MAP_READONLY in
+           OMP_CLAUSE_MAP
+
+       OMP_CLAUSE__CACHE__READONLY in
+           OMP_CLAUSE__CACHE_
+
    constant_flag:
 
        TREE_CONSTANT in

diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc
index 654f5247e3acf8fc9ebff986e5ee2aea9cf35f46..926f7e006a7618db1c4c8cdcea8282c831f6bf54 100644 (file)
--- 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)
 
     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, 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 e1fc6c2221ddaf87daaacf93438acc7c791f32cd..b67a37d6522662b89ea8e8170ab10c56ecaffb95 100644 (file)
--- 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_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) \