tree-core.h (enum omp_clause_code): Adjust OMP_CLAUSE_USE_DEVICE_PTR OpenMP description.

	* tree-core.h (enum omp_clause_code): Adjust OMP_CLAUSE_USE_DEVICE_PTR
	OpenMP description.  Add OMP_CLAUSE_USE_DEVICE_ADDR clause.
	* tree.c (omp_clause_num_ops, omp_clause_code_name): Add entries
	for OMP_CLAUSE_USE_DEVICE_ADDR clause.
	(walk_tree_1): Handle OMP_CLAUSE_USE_DEVICE_ADDR.
	* tree-pretty-print.c (dump_omp_clause): Likewise.
	* tree-nested.c (convert_nonlocal_omp_clauses,
	convert_local_omp_clauses): Likewise.
	* gimplify.c (gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses):
	Likewise.
	* omp-low.c (scan_sharing_clauses, lower_omp_target): Likewise.
	Treat OMP_CLAUSE_USE_DEVICE_ADDR like OMP_CLAUSE_USE_DEVICE_PTR
	clause with array or reference to array types, no matter what type
	except for reference it has.
gcc/c-family/
	* c-pragma.h (enum pragma_omp_clause): Add
	PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR.  Set PRAGMA_OACC_CLAUSE_USE_DEVICE
	equal to PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR instead of being a separate
	enumeration value.
gcc/c/
	* c-parser.c (c_parser_omp_clause_name): Parse use_device_addr clause.
	(c_parser_omp_clause_use_device_addr): New function.
	(c_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR.
	(OMP_TARGET_DATA_CLAUSE_MASK): Add PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR.
	(c_parser_omp_target_data): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR
	like PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR, adjust diagnostics about no
	map or use_device_* clauses.
	* c-typeck.c (c_finish_omp_clauses): For OMP_CLAUSE_USE_DEVICE_PTR
	in OpenMP, require pointer type rather than pointer or array type.
	Handle OMP_CLAUSE_USE_DEVICE_ADDR.
gcc/cp/
	* parser.c (cp_parser_omp_clause_name): Parse use_device_addr clause.
	(cp_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR.
	(OMP_TARGET_DATA_CLAUSE_MASK): Add PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR.
	(cp_parser_omp_target_data): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR
	like PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR, adjust diagnostics about no
	map or use_device_* clauses.
	* semantics.c (finish_omp_clauses): For OMP_CLAUSE_USE_DEVICE_PTR
	in OpenMP, require pointer or reference to pointer type rather than
	pointer or array or reference to pointer or array type. Handle
	OMP_CLAUSE_USE_DEVICE_ADDR.
	* pt.c (tsubst_omp_clauses): Handle OMP_CLAUSE_USE_DEVICE_ADDR.
gcc/testsuite/
	* c-c++-common/gomp/target-data-1.c (foo): Use use_device_addr clause
	instead of use_device_ptr clause where required by OpenMP 5.0, add
	further tests for both use_device_ptr and use_device_addr clauses.
libgomp/
	* testsuite/libgomp.c/target-18.c (struct S): New type.
	(foo): Use use_device_addr clause instead of use_device_ptr clause
	where required by OpenMP 5.0, add further tests for both use_device_ptr
	and use_device_addr clauses.
	* testsuite/libgomp.c++/target-9.C (struct S): New type.
	(foo): Use use_device_addr clause instead of use_device_ptr clause
	where required by OpenMP 5.0, add further tests for both use_device_ptr
	and use_device_addr clauses.  Add t and u arguments.
	(main): Adjust caller.

From-SVN: r274159

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index b85a1dba1052b89e54b1073fbfd91d9f6aee728d..23837ab3a9a097e92ff10bc68cd6473cb2f2b4d7 100644 (file)
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,20 @@
+2019-08-07  Jakub Jelinek  <jakub@redhat.com>
+
+       * tree-core.h (enum omp_clause_code): Adjust OMP_CLAUSE_USE_DEVICE_PTR
+       OpenMP description.  Add OMP_CLAUSE_USE_DEVICE_ADDR clause.
+       * tree.c (omp_clause_num_ops, omp_clause_code_name): Add entries
+       for OMP_CLAUSE_USE_DEVICE_ADDR clause.
+       (walk_tree_1): Handle OMP_CLAUSE_USE_DEVICE_ADDR.
+       * tree-pretty-print.c (dump_omp_clause): Likewise.
+       * tree-nested.c (convert_nonlocal_omp_clauses,
+       convert_local_omp_clauses): Likewise.
+       * gimplify.c (gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses):
+       Likewise.
+       * omp-low.c (scan_sharing_clauses, lower_omp_target): Likewise.
+       Treat OMP_CLAUSE_USE_DEVICE_ADDR like OMP_CLAUSE_USE_DEVICE_PTR
+       clause with array or reference to array types, no matter what type
+       except for reference it has.
+
 2019-08-07  Kewen Lin  <linkw@gcc.gnu.org>
 
        * config/rs6000/vector.md (vrotr<mode>3): New define_expand.

diff --git a/gcc/c-family/ChangeLog b/gcc/c-family/ChangeLog
index 693a546347c0ea33be27fe57ce4e12c4580d0948..873b9ad5b7a3a96926a375cac9b9f6c5b2add11e 100644 (file)
--- a/gcc/c-family/ChangeLog
+++ b/gcc/c-family/ChangeLog
@@ -1,3 +1,10 @@
+2019-08-07  Jakub Jelinek  <jakub@redhat.com>
+
+       * c-pragma.h (enum pragma_omp_clause): Add
+       PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR.  Set PRAGMA_OACC_CLAUSE_USE_DEVICE
+       equal to PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR instead of being a separate
+       enumeration value.
+
 2019-08-05  Marek Polacek  <polacek@redhat.com>
 
        PR c++/91338 - Implement P1161R3: Deprecate a[b,c].

diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index e8a509f70733fc7af7e21a6919eb6430e7f3b85d..803dc1b2ac63040840a538bb8a69686cdb7665ea 100644 (file)
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -137,6 +137,7 @@ enum pragma_omp_clause {
   PRAGMA_OMP_CLAUSE_UNIFORM,
   PRAGMA_OMP_CLAUSE_UNTIED,
   PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR,
+  PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR,
 
   /* Clauses for OpenACC.  */
   PRAGMA_OACC_CLAUSE_ASYNC,
@@ -157,7 +158,6 @@ enum pragma_omp_clause {
   PRAGMA_OACC_CLAUSE_SELF,
   PRAGMA_OACC_CLAUSE_SEQ,
   PRAGMA_OACC_CLAUSE_TILE,
-  PRAGMA_OACC_CLAUSE_USE_DEVICE,
   PRAGMA_OACC_CLAUSE_VECTOR,
   PRAGMA_OACC_CLAUSE_VECTOR_LENGTH,
   PRAGMA_OACC_CLAUSE_WAIT,
@@ -171,7 +171,8 @@ enum pragma_omp_clause {
   PRAGMA_OACC_CLAUSE_IF = PRAGMA_OMP_CLAUSE_IF,
   PRAGMA_OACC_CLAUSE_PRIVATE = PRAGMA_OMP_CLAUSE_PRIVATE,
   PRAGMA_OACC_CLAUSE_REDUCTION = PRAGMA_OMP_CLAUSE_REDUCTION,
-  PRAGMA_OACC_CLAUSE_LINK = PRAGMA_OMP_CLAUSE_LINK
+  PRAGMA_OACC_CLAUSE_LINK = PRAGMA_OMP_CLAUSE_LINK,
+  PRAGMA_OACC_CLAUSE_USE_DEVICE = PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR
 };
 
 extern struct cpp_reader* parse_in;

diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog
index 18c4e413e2d38ef5c55400738ba1acdc120c5041..18c1ce0b6df7d2b0e679461eb2619dea9b5d9475 100644 (file)
--- a/gcc/c/ChangeLog
+++ b/gcc/c/ChangeLog
@@ -1,3 +1,16 @@
+2019-08-07  Jakub Jelinek  <jakub@redhat.com>
+
+       * c-parser.c (c_parser_omp_clause_name): Parse use_device_addr clause.
+       (c_parser_omp_clause_use_device_addr): New function.
+       (c_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR.
+       (OMP_TARGET_DATA_CLAUSE_MASK): Add PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR.
+       (c_parser_omp_target_data): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR
+       like PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR, adjust diagnostics about no
+       map or use_device_* clauses.
+       * c-typeck.c (c_finish_omp_clauses): For OMP_CLAUSE_USE_DEVICE_PTR
+       in OpenMP, require pointer type rather than pointer or array type.
+       Handle OMP_CLAUSE_USE_DEVICE_ADDR.
+
 2019-07-31  Jakub Jelinek  <jakub@redhat.com>
 
        PR c/91192

diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index c7c13a6d9ae3c2fa7f12ae0d6ecbd6244cd9f189..9b5cfcf50d949d1ecfb62db4076a43053ddb82e8 100644 (file)
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -11866,6 +11866,8 @@ c_parser_omp_clause_name (c_parser *parser)
            result = PRAGMA_OMP_CLAUSE_UNTIED;
          else if (!strcmp ("use_device", p))
            result = PRAGMA_OACC_CLAUSE_USE_DEVICE;
+         else if (!strcmp ("use_device_addr", p))
+           result = PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR;
          else if (!strcmp ("use_device_ptr", p))
            result = PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR;
          break;
@@ -13121,6 +13123,16 @@ c_parser_omp_clause_use_device_ptr (c_parser *parser, tree list)
                                       list);
 }
 
+/* OpenMP 5.0:
+   use_device_addr ( variable-list ) */
+
+static tree
+c_parser_omp_clause_use_device_addr (c_parser *parser, tree list)
+{
+  return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_USE_DEVICE_ADDR,
+                                      list);
+}
+
 /* OpenMP 4.5:
    is_device_ptr ( variable-list ) */
 
@@ -15321,6 +15333,10 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask,
          clauses = c_parser_omp_clause_use_device_ptr (parser, clauses);
          c_name = "use_device_ptr";
          break;
+       case PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR:
+         clauses = c_parser_omp_clause_use_device_addr (parser, clauses);
+         c_name = "use_device_addr";
+         break;
        case PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR:
          clauses = c_parser_omp_clause_is_device_ptr (parser, clauses);
          c_name = "is_device_ptr";
@@ -18288,7 +18304,8 @@ c_parser_omp_teams (location_t loc, c_parser *parser,
        ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE)       \
        | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_MAP)          \
        | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF)           \
-       | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR))
+       | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR) \
+       | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR))
 
 static tree
 c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p)
@@ -18323,7 +18340,8 @@ c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p)
            *pc = OMP_CLAUSE_CHAIN (*pc);
            continue;
          }
-      else if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_USE_DEVICE_PTR)
+      else if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_USE_DEVICE_PTR
+              || OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_USE_DEVICE_ADDR)
        map_seen = 3;
       pc = &OMP_CLAUSE_CHAIN (*pc);
     }
@@ -18333,7 +18351,8 @@ c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p)
       if (map_seen == 0)
        error_at (loc,
                  "%<#pragma omp target data%> must contain at least "
-                 "one %<map%> or %<use_device_ptr%> clause");
+                 "one %<map%>, %<use_device_ptr%> or %<use_device_addr%> "
+                 "clause");
       return NULL_TREE;
     }
 

diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index 9a1a9106a40bd008549fe57f09b832bca05077e2..bda90fdd27350555ef9df6df3ee5a2d1110c82ea 100644 (file)
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -14609,16 +14609,32 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
        case OMP_CLAUSE_IS_DEVICE_PTR:
        case OMP_CLAUSE_USE_DEVICE_PTR:
          t = OMP_CLAUSE_DECL (c);
-         if (TREE_CODE (TREE_TYPE (t)) != POINTER_TYPE
-             && TREE_CODE (TREE_TYPE (t)) != ARRAY_TYPE)
+         if (TREE_CODE (TREE_TYPE (t)) != POINTER_TYPE)
            {
-             error_at (OMP_CLAUSE_LOCATION (c),
-                       "%qs variable is neither a pointer nor an array",
-                       omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
-             remove = true;
+             if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR
+                 && ort == C_ORT_OMP)
+               {
+                 error_at (OMP_CLAUSE_LOCATION (c),
+                           "%qs variable is not a pointer",
+                           omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+                 remove = true;
+               }
+             else if (TREE_CODE (TREE_TYPE (t)) != ARRAY_TYPE)
+               {
+                 error_at (OMP_CLAUSE_LOCATION (c),
+                           "%qs variable is neither a pointer nor an array",
+                           omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+                 remove = true;
+               }
            }
          goto check_dup_generic;
 
+       case OMP_CLAUSE_USE_DEVICE_ADDR:
+         t = OMP_CLAUSE_DECL (c);
+         if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
+           c_mark_addressable (t);
+         goto check_dup_generic;
+
        case OMP_CLAUSE_NOWAIT:
          if (copyprivate_seen)
            {

diff --git a/gcc/cp/ChangeLog b/gcc/cp/ChangeLog
index 9b0a34c98aac449d5a89cfd29b3a49a9f70fc8a7..c1e977b6926318928582495faafbc12c25d1dafc 100644 (file)
--- a/gcc/cp/ChangeLog
+++ b/gcc/cp/ChangeLog
@@ -1,3 +1,17 @@
+2019-08-07  Jakub Jelinek  <jakub@redhat.com>
+
+       * parser.c (cp_parser_omp_clause_name): Parse use_device_addr clause.
+       (cp_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR.
+       (OMP_TARGET_DATA_CLAUSE_MASK): Add PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR.
+       (cp_parser_omp_target_data): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR
+       like PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR, adjust diagnostics about no
+       map or use_device_* clauses.
+       * semantics.c (finish_omp_clauses): For OMP_CLAUSE_USE_DEVICE_PTR
+       in OpenMP, require pointer or reference to pointer type rather than
+       pointer or array or reference to pointer or array type. Handle
+       OMP_CLAUSE_USE_DEVICE_ADDR.
+       * pt.c (tsubst_omp_clauses): Handle OMP_CLAUSE_USE_DEVICE_ADDR.
+
 2019-08-06  Jason Merrill  <jason@redhat.com>
 
        PR c++/91378 - ICE with noexcept and auto return type.

diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 79da7b52eb9a27558bb0b8777ca3bdd5216bc367..4d07a6a301126426c4b0eb9ee32bd1b89055100a 100644 (file)
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -32648,6 +32648,8 @@ cp_parser_omp_clause_name (cp_parser *parser)
            result = PRAGMA_OMP_CLAUSE_UNTIED;
          else if (!strcmp ("use_device", p))
            result = PRAGMA_OACC_CLAUSE_USE_DEVICE;
+         else if (!strcmp ("use_device_addr", p))
+           result = PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR;
          else if (!strcmp ("use_device_ptr", p))
            result = PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR;
          break;
@@ -35637,6 +35639,11 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask,
                                            clauses);
          c_name = "use_device_ptr";
          break;
+       case PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR:
+         clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_USE_DEVICE_ADDR,
+                                           clauses);
+         c_name = "use_device_addr";
+         break;
        case PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR:
          clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_IS_DEVICE_PTR,
                                            clauses);
@@ -38715,7 +38722,8 @@ cp_parser_omp_teams (cp_parser *parser, cp_token *pragma_tok,
        ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE)       \
        | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_MAP)          \
        | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF)           \
-       | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR))
+       | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR) \
+       | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR))
 
 static tree
 cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p)
@@ -38751,7 +38759,8 @@ cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p)
            *pc = OMP_CLAUSE_CHAIN (*pc);
            continue;
          }
-      else if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_USE_DEVICE_PTR)
+      else if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_USE_DEVICE_PTR
+              || OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_USE_DEVICE_ADDR)
        map_seen = 3;
       pc = &OMP_CLAUSE_CHAIN (*pc);
     }
@@ -38761,7 +38770,8 @@ cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p)
       if (map_seen == 0)
        error_at (pragma_tok->location,
                  "%<#pragma omp target data%> must contain at least "
-                 "one %<map%> or %<use_device_ptr%> clause");
+                 "one %<map%>, %<use_device_ptr%> or %<use_device_addr%> "
+                 "clause");
       return NULL_TREE;
     }
 

diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c
index b71fbaad78991b7eb6674162bdeab5f061fea944..b1ad99d148199670056a1ebe76da31b38b6d90dc 100644 (file)
--- a/gcc/cp/pt.c
+++ b/gcc/cp/pt.c
@@ -16303,6 +16303,7 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort,
        case OMP_CLAUSE_MAP:
        case OMP_CLAUSE_NONTEMPORAL:
        case OMP_CLAUSE_USE_DEVICE_PTR:
+       case OMP_CLAUSE_USE_DEVICE_ADDR:
        case OMP_CLAUSE_IS_DEVICE_PTR:
        case OMP_CLAUSE_INCLUSIVE:
        case OMP_CLAUSE_EXCLUSIVE:
@@ -16427,6 +16428,7 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort,
          case OMP_CLAUSE_IN_REDUCTION:
          case OMP_CLAUSE_TASK_REDUCTION:
          case OMP_CLAUSE_USE_DEVICE_PTR:
+         case OMP_CLAUSE_USE_DEVICE_ADDR:
          case OMP_CLAUSE_IS_DEVICE_PTR:
          case OMP_CLAUSE_INCLUSIVE:
          case OMP_CLAUSE_EXCLUSIVE:

diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index fa6962454bf4357980eccf056f813953b8689d59..77e7a6dced2d0484f1438d23e7bf24f000ac71c8 100644 (file)
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -7524,20 +7524,41 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
            {
              tree type = TREE_TYPE (t);
              if (!TYPE_PTR_P (type)
-                 && TREE_CODE (type) != ARRAY_TYPE
-                 && (!TYPE_REF_P (type)
-                     || (!TYPE_PTR_P (TREE_TYPE (type))
-                         && TREE_CODE (TREE_TYPE (type)) != ARRAY_TYPE)))
+                 && (!TYPE_REF_P (type) || !TYPE_PTR_P (TREE_TYPE (type))))
                {
-                 error_at (OMP_CLAUSE_LOCATION (c),
-                           "%qs variable is neither a pointer, nor an array "
-                           "nor reference to pointer or array",
-                           omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
-                 remove = true;
+                 if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR
+                     && ort == C_ORT_OMP)
+                   {
+                     error_at (OMP_CLAUSE_LOCATION (c),
+                               "%qs variable is neither a pointer "
+                               "nor reference to pointer",
+                               omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+                     remove = true;
+                   }
+                 else if (TREE_CODE (type) != ARRAY_TYPE
+                          && (!TYPE_REF_P (type)
+                              || TREE_CODE (TREE_TYPE (type)) != ARRAY_TYPE))
+                   {
+                     error_at (OMP_CLAUSE_LOCATION (c),
+                               "%qs variable is neither a pointer, nor an "
+                               "array nor reference to pointer or array",
+                               omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+                     remove = true;
+                   }
                }
            }
          goto check_dup_generic;
 
+       case OMP_CLAUSE_USE_DEVICE_ADDR:
+         field_ok = true;
+         t = OMP_CLAUSE_DECL (c);
+         if (!processing_template_decl
+             && (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
+             && !TYPE_REF_P (TREE_TYPE (t))
+             && !cxx_mark_addressable (t))
+           remove = true;
+         goto check_dup_generic;
+
        case OMP_CLAUSE_NOWAIT:
        case OMP_CLAUSE_DEFAULT:
        case OMP_CLAUSE_UNTIED:

diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 10b9b68d5ce633491fd9ad83d5813a5abc5568fe..978df86f72a608b6f3d8f32b4cb7092c08238018 100644 (file)
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -9015,8 +9015,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
          goto do_notice;
 
        case OMP_CLAUSE_USE_DEVICE_PTR:
-         flags = GOVD_FIRSTPRIVATE | GOVD_EXPLICIT;
-         goto do_add;
+       case OMP_CLAUSE_USE_DEVICE_ADDR:
        case OMP_CLAUSE_IS_DEVICE_PTR:
          flags = GOVD_FIRSTPRIVATE | GOVD_EXPLICIT;
          goto do_add;
@@ -10264,6 +10263,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
        case OMP_CLAUSE_ORDER:
        case OMP_CLAUSE_BIND:
        case OMP_CLAUSE_USE_DEVICE_PTR:
+       case OMP_CLAUSE_USE_DEVICE_ADDR:
        case OMP_CLAUSE_IS_DEVICE_PTR:
        case OMP_CLAUSE_ASYNC:
        case OMP_CLAUSE_WAIT:

diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 4a6ea0a1b71586411b0b96abbf548c0cd3d795cc..ca8edde6cdbc4a4a96600aeae2ac67bb3b5521bd 100644 (file)
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -1238,8 +1238,11 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
          break;
 
        case OMP_CLAUSE_USE_DEVICE_PTR:
+       case OMP_CLAUSE_USE_DEVICE_ADDR:
          decl = OMP_CLAUSE_DECL (c);
-         if (TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
+         if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR
+              && !omp_is_reference (decl))
+             || TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
            install_var_field (decl, true, 3, ctx);
          else
            install_var_field (decl, false, 3, ctx);
@@ -1635,6 +1638,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
        case OMP_CLAUSE_ORDER:
        case OMP_CLAUSE_BIND:
        case OMP_CLAUSE_USE_DEVICE_PTR:
+       case OMP_CLAUSE_USE_DEVICE_ADDR:
        case OMP_CLAUSE_NONTEMPORAL:
        case OMP_CLAUSE_ASYNC:
        case OMP_CLAUSE_WAIT:
@@ -11465,6 +11469,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
        break;
 
       case OMP_CLAUSE_USE_DEVICE_PTR:
+      case OMP_CLAUSE_USE_DEVICE_ADDR:
       case OMP_CLAUSE_IS_DEVICE_PTR:
        var = OMP_CLAUSE_DECL (c);
        map_cnt++;
@@ -11481,7 +11486,9 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
            SET_DECL_VALUE_EXPR (new_var, x);
            DECL_HAS_VALUE_EXPR_P (new_var) = 1;
          }
-       else if (TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
+       else if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR
+                 && !omp_is_reference (var))
+                || TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
          {
            tree new_var = lookup_decl (var, ctx);
            tree type = build_pointer_type (TREE_TYPE (var));
@@ -11846,23 +11853,27 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
            break;
 
          case OMP_CLAUSE_USE_DEVICE_PTR:
+         case OMP_CLAUSE_USE_DEVICE_ADDR:
          case OMP_CLAUSE_IS_DEVICE_PTR:
            ovar = OMP_CLAUSE_DECL (c);
            var = lookup_decl_in_outer_ctx (ovar, ctx);
            x = build_sender_ref (ovar, ctx);
-           if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR)
+           if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR)
              tkind = GOMP_MAP_USE_DEVICE_PTR;
            else
              tkind = GOMP_MAP_FIRSTPRIVATE_INT;
            type = TREE_TYPE (ovar);
-           if (TREE_CODE (type) == ARRAY_TYPE)
+           if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR
+                && !omp_is_reference (ovar))
+               || TREE_CODE (type) == ARRAY_TYPE)
              var = build_fold_addr_expr (var);
            else
              {
                if (omp_is_reference (ovar))
                  {
                    type = TREE_TYPE (type);
-                   if (TREE_CODE (type) != ARRAY_TYPE)
+                   if (TREE_CODE (type) != ARRAY_TYPE
+                       && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_USE_DEVICE_ADDR)
                      var = build_simple_mem_ref (var);
                    var = fold_convert (TREE_TYPE (x), var);
                  }
@@ -12017,9 +12028,10 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
              }
            break;
          case OMP_CLAUSE_USE_DEVICE_PTR:
+         case OMP_CLAUSE_USE_DEVICE_ADDR:
          case OMP_CLAUSE_IS_DEVICE_PTR:
            var = OMP_CLAUSE_DECL (c);
-           if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR)
+           if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR)
              x = build_sender_ref (var, ctx);
            else
              x = build_receiver_ref (var, false, ctx);
@@ -12034,7 +12046,9 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
                gimple_seq_add_stmt (&new_body,
                                     gimple_build_assign (new_var, x));
              }
-           else if (TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
+           else if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR
+                     && !omp_is_reference (var))
+                    || TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
              {
                tree new_var = lookup_decl (var, ctx);
                new_var = DECL_VALUE_EXPR (new_var);
@@ -12052,7 +12066,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
                if (omp_is_reference (var))
                  {
                    type = TREE_TYPE (type);
-                   if (TREE_CODE (type) != ARRAY_TYPE)
+                   if (TREE_CODE (type) != ARRAY_TYPE
+                       && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_USE_DEVICE_ADDR)
                      {
                        tree v = create_tmp_var_raw (type, get_name (var));
                        gimple_add_tmp_var (v);

diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index a372d7fbcf17caf4fb9727b98df103710bb8a06a..d524f72b4d299e0ec57d35f80df6a29dcb297f31 100644 (file)
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,3 +1,9 @@
+2019-08-07  Jakub Jelinek  <jakub@redhat.com>
+
+       * c-c++-common/gomp/target-data-1.c (foo): Use use_device_addr clause
+       instead of use_device_ptr clause where required by OpenMP 5.0, add
+       further tests for both use_device_ptr and use_device_addr clauses.
+
 2019-08-07  Kewen Lin  <linkw@gcc.gnu.org>
 
        * gcc.target/powerpc/vec_rotate-1.c: New test.

diff --git a/gcc/testsuite/c-c++-common/gomp/target-data-1.c b/gcc/testsuite/c-c++-common/gomp/target-data-1.c
index 0d4975bcc03d8296a969b1de690ad53276ba8695..7aa111a92a46c9046942906ecdfd83a5e3fc460b 100644 (file)
--- a/gcc/testsuite/c-c++-common/gomp/target-data-1.c
+++ b/gcc/testsuite/c-c++-common/gomp/target-data-1.c
@@ -4,15 +4,39 @@ void
 foo (void)
 {
   int a[4] = { 1, 2, 3, 4 };
+  int *p = &a[0];
+  int x = 5;
+  #pragma omp target data map(to:p[:4])
+  #pragma omp target data use_device_ptr(p)
+  #pragma omp target is_device_ptr(p)
+  {
+    p[0]++;
+  }
   #pragma omp target data map(to:a)
-  #pragma omp target data use_device_ptr(a)
+  #pragma omp target data use_device_addr(a)
   #pragma omp target is_device_ptr(a)
   {
-    a[0]++;
+    p[0]++;
+  }
+  #pragma omp target data map(to:x)
+  #pragma omp target data use_device_addr(x)
+  {
+    int *q = &x;
+    #pragma omp target is_device_ptr(q)
+    {
+      q[0]++;
+    }
   }
   #pragma omp target data              /* { dg-error "must contain at least one" } */
   a[0]++;
+  #pragma omp target data map(to:p)
+  #pragma omp target data use_device_ptr(p) use_device_ptr(p) /* { dg-error "appears more than once in data clauses" } */
+  a[0]++;
   #pragma omp target data map(to:a)
-  #pragma omp target data use_device_ptr(a) use_device_ptr(a) /* { dg-error "appears more than once in data clauses" } */
+  #pragma omp target data use_device_addr(a) use_device_addr(a) /* { dg-error "appears more than once in data clauses" } */
   a[0]++;
+  #pragma omp target data map(to:a)
+  #pragma omp target data use_device_ptr(a)    /* { dg-error "'use_device_ptr' variable is not a pointer" "" { target c } } */
+                                               /* { dg-error "'use_device_ptr' variable is neither a pointer nor reference to pointer" "" { target c++ } .-1 } */
+  a[0]++;                                      /* { dg-error "must contain at least one" "" { target *-*-* } .-2 } */
 }

diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index 471956269121f96fbe27d72ff75876e8d65b298d..8216a60671803207a8ab2a663251afe419103922 100644 (file)
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -307,9 +307,12 @@ enum omp_clause_code {
   OMP_CLAUSE_MAP,
 
   /* OpenACC clause: use_device (variable-list).
-     OpenMP clause: use_device_ptr (variable-list).  */
+     OpenMP clause: use_device_ptr (ptr-list).  */
   OMP_CLAUSE_USE_DEVICE_PTR,
 
+  /* OpenMP clause: use_device_addr (variable-list).  */
+  OMP_CLAUSE_USE_DEVICE_ADDR,
+
   /* OpenMP clause: is_device_ptr (variable-list).  */
   OMP_CLAUSE_IS_DEVICE_PTR,
 

diff --git a/gcc/tree-nested.c b/gcc/tree-nested.c
index 5d189572e19160cb232f658e92cd325e4ae91cca..4df07edcf0a8bb31c33d71f7442aa1e343a14727 100644 (file)
--- a/gcc/tree-nested.c
+++ b/gcc/tree-nested.c
@@ -1227,6 +1227,7 @@ convert_nonlocal_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
        case OMP_CLAUSE_TO_DECLARE:
        case OMP_CLAUSE_LINK:
        case OMP_CLAUSE_USE_DEVICE_PTR:
+       case OMP_CLAUSE_USE_DEVICE_ADDR:
        case OMP_CLAUSE_IS_DEVICE_PTR:
        do_decl_clause:
          decl = OMP_CLAUSE_DECL (clause);
@@ -1947,6 +1948,7 @@ convert_local_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
        case OMP_CLAUSE_TO_DECLARE:
        case OMP_CLAUSE_LINK:
        case OMP_CLAUSE_USE_DEVICE_PTR:
+       case OMP_CLAUSE_USE_DEVICE_ADDR:
        case OMP_CLAUSE_IS_DEVICE_PTR:
        do_decl_clause:
          decl = OMP_CLAUSE_DECL (clause);

diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index 9bea132f7f02a762f6045deb575d8df941077a97..8091a0af7f34baa678b96f8c3a6197e2278ad42d 100644 (file)
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -465,6 +465,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
     case OMP_CLAUSE_USE_DEVICE_PTR:
       name = "use_device_ptr";
       goto print_remap;
+    case OMP_CLAUSE_USE_DEVICE_ADDR:
+      name = "use_device_addr";
+      goto print_remap;
     case OMP_CLAUSE_IS_DEVICE_PTR:
       name = "is_device_ptr";
       goto print_remap;

diff --git a/gcc/tree.c b/gcc/tree.c
index efa49e99d659a23ffc55fb31c67a793003e0bdde..1f2aac4ffa554ac9c1d3f42bc02e4b928360e654 100644 (file)
--- a/gcc/tree.c
+++ b/gcc/tree.c
@@ -299,6 +299,7 @@ unsigned const char omp_clause_num_ops[] =
   2, /* OMP_CLAUSE_TO  */
   2, /* OMP_CLAUSE_MAP  */
   1, /* OMP_CLAUSE_USE_DEVICE_PTR  */
+  1, /* OMP_CLAUSE_USE_DEVICE_ADDR  */
   1, /* OMP_CLAUSE_IS_DEVICE_PTR  */
   1, /* OMP_CLAUSE_INCLUSIVE  */
   1, /* OMP_CLAUSE_EXCLUSIVE  */
@@ -382,6 +383,7 @@ const char * const omp_clause_code_name[] =
   "to",
   "map",
   "use_device_ptr",
+  "use_device_addr",
   "is_device_ptr",
   "inclusive",
   "exclusive",
@@ -12384,6 +12386,7 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,
        case OMP_CLAUSE_TO_DECLARE:
        case OMP_CLAUSE_LINK:
        case OMP_CLAUSE_USE_DEVICE_PTR:
+       case OMP_CLAUSE_USE_DEVICE_ADDR:
        case OMP_CLAUSE_IS_DEVICE_PTR:
        case OMP_CLAUSE_INCLUSIVE:
        case OMP_CLAUSE_EXCLUSIVE:

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 0331c51c6394e2348282927c791050ee699742ba..01440f7a2521543cbe8b3e6ed7e68e189eb31963 100644 (file)
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,15 @@
+2019-08-07  Jakub Jelinek  <jakub@redhat.com>
+
+       * testsuite/libgomp.c/target-18.c (struct S): New type.
+       (foo): Use use_device_addr clause instead of use_device_ptr clause
+       where required by OpenMP 5.0, add further tests for both use_device_ptr
+       and use_device_addr clauses.
+       * testsuite/libgomp.c++/target-9.C (struct S): New type.
+       (foo): Use use_device_addr clause instead of use_device_ptr clause
+       where required by OpenMP 5.0, add further tests for both use_device_ptr
+       and use_device_addr clauses.  Add t and u arguments.
+       (main): Adjust caller.
+
 2019-08-06  Jakub Jelinek  <jakub@redhat.com>
 
        * testsuite/libgomp.c++/loop-13.C: New test.

diff --git a/libgomp/testsuite/libgomp.c++/target-9.C b/libgomp/testsuite/libgomp.c++/target-9.C
index a5d171b0b3d897375fcaefe272739e7f26e7bef5..83a61cf599b45933ffb95062223fe08136ab98c3 100644 (file)
--- a/libgomp/testsuite/libgomp.c++/target-9.C
+++ b/libgomp/testsuite/libgomp.c++/target-9.C
@@ -1,10 +1,13 @@
 extern "C" void abort (void);
+struct S { int e, f; };
 
 void
-foo (int *&p, int (&s)[5], int n)
+foo (int *&p, int (&s)[5], int &t, S &u, int n)
 {
   int a[4] = { 7, 8, 9, 10 }, b[n], c[3] = { 20, 21, 22 };
   int *r = a + 1, *q = p - 1, i, err;
+  int v = 27;
+  S w = { 28, 29 };
   for (i = 0; i < n; i++)
     b[i] = 9 + i;
   #pragma omp target data map(to:a)
@@ -30,7 +33,7 @@ foo (int *&p, int (&s)[5], int n)
   if (err)
     abort ();
   #pragma omp target data map(to:b)
-  #pragma omp target data use_device_ptr(b) map(from:err)
+  #pragma omp target data use_device_addr(b) map(from:err)
   #pragma omp target is_device_ptr(b) private(i) map(from:err)
   {
     err = 0;
@@ -41,7 +44,7 @@ foo (int *&p, int (&s)[5], int n)
   if (err)
     abort ();
   #pragma omp target data map(to:c)
-  #pragma omp target data use_device_ptr(c) map(from:err)
+  #pragma omp target data use_device_addr(c) map(from:err)
   #pragma omp target is_device_ptr(c) private(i) map(from:err)
   {
     err = 0;
@@ -52,7 +55,7 @@ foo (int *&p, int (&s)[5], int n)
   if (err)
     abort ();
   #pragma omp target data map(to:s[:5])
-  #pragma omp target data use_device_ptr(s) map(from:err)
+  #pragma omp target data use_device_addr(s) map(from:err)
   #pragma omp target is_device_ptr(s) private(i) map(from:err)
   {
     err = 0;
@@ -62,6 +65,34 @@ foo (int *&p, int (&s)[5], int n)
   }
   if (err)
     abort ();
+  #pragma omp target data map(to: v) map(to:u)
+  #pragma omp target data use_device_addr (v) use_device_addr (u) map(from:err)
+  {
+    int *z = &v;
+    S *x = &u;
+    #pragma omp target is_device_ptr (z, x) map(from:err)
+    {
+      err = 0;
+      if (*z != 27 || x->e != 25 || x->f != 26)
+       err = 1;
+    }
+  }
+  if (err)
+    abort ();
+  #pragma omp target data map(to: t, w)
+  #pragma omp target data use_device_addr (t, w) map(from:err)
+  {
+    int *z = &t;
+    S *x = &w;
+    #pragma omp target is_device_ptr (z) is_device_ptr (x) map(from:err)
+    {
+      err = 0;
+      if (*z != 24 || x->e != 28 || x->f != 29)
+       err = 1;
+    }
+  }
+  if (err)
+    abort ();
 }
 
 int
@@ -69,5 +100,7 @@ main ()
 {
   int a[4] = { 0, 1, 2, 3 }, b[5] = { 17, 18, 19, 20, 21 };
   int *p = a + 1;
-  foo (p, b, 9);
+  int t = 24;
+  S u = { 25, 26 };
+  foo (p, b, t, u, 9);
 }

diff --git a/libgomp/testsuite/libgomp.c/target-18.c b/libgomp/testsuite/libgomp.c/target-18.c
index cbacaf6a77ac673864f943a88e312e743441155f..dd511fb98cb8fee00e77351797aed57e9bd8354a 100644 (file)
--- a/libgomp/testsuite/libgomp.c/target-18.c
+++ b/libgomp/testsuite/libgomp.c/target-18.c
@@ -1,9 +1,11 @@
 extern void abort (void);
+struct S { int e, f; };
 
 void
 foo (int n)
 {
-  int a[4] = { 0, 1, 2, 3 }, b[n];
+  int a[4] = { 0, 1, 2, 3 }, b[n], c = 4;
+  struct S d = { 5, 6 };
   int *p = a + 1, i, err;
   for (i = 0; i < n; i++)
     b[i] = 9 + i;
@@ -21,7 +23,7 @@ foo (int n)
   for (i = 0; i < 4; i++)
     a[i] = 23 + i;
   #pragma omp target data map(to:a)
-  #pragma omp target data use_device_ptr(a) map(from:err)
+  #pragma omp target data use_device_addr(a) map(from:err)
   #pragma omp target is_device_ptr(a) private(i) map(from:err)
   {
     err = 0;
@@ -32,7 +34,7 @@ foo (int n)
   if (err)
     abort ();
   #pragma omp target data map(to:b)
-  #pragma omp target data use_device_ptr(b) map(from:err)
+  #pragma omp target data use_device_addr(b) map(from:err)
   #pragma omp target is_device_ptr(b) private(i) map(from:err)
   {
     err = 0;
@@ -42,6 +44,28 @@ foo (int n)
   }
   if (err)
     abort ();
+  #pragma omp target data map(to:c)
+  #pragma omp target data use_device_addr(c) map(from:err)
+  {
+    int *q = &c;
+    #pragma omp target is_device_ptr(q) map(from:err)
+    {
+      err = *q != 4;
+    }
+  }
+  if (err)
+    abort ();
+  #pragma omp target data map(to:d)
+  #pragma omp target data use_device_addr(d) map(from:err)
+  {
+    struct S *r = &d;
+    #pragma omp target is_device_ptr(r) map(from:err)
+    {
+      err = r->e != 5 || r->f != 6;
+    }
+  }
+  if (err)
+    abort ();
 }
 
 int