OpenMP: Add parser support for target's device_type clause

Enables the existing 'device_type(nohost|host|any)' for the
'target' construct; for now, it will fail with a 'sorry,
unimplemented' for all but 'any'.

gcc/c-family/ChangeLog:

	* c-omp.cc (c_omp_split_clauses): Handle target's
	device_type clause.

gcc/c/ChangeLog:

	* c-parser.cc (OMP_TARGET_CLAUSE_MASK): Add
	device_type clause.

gcc/cp/ChangeLog:

	* parser.cc (OMP_TARGET_CLAUSE_MASK): Add
	device_type clause.

gcc/fortran/ChangeLog:

	* dump-parse-tree.cc (show_omp_clauses): Handle
	device_type clause.
	* openmp.cc (gfc_match_omp_clauses): Reorder to
	match 'device' after 'device_...' to avoid parse
	errors.
	(OMP_TARGET_CLAUSES): Add device_type clause.
	* trans-openmp.cc (gfc_trans_omp_clauses,
	gfc_split_omp_clauses): Handle device_type clause.

gcc/ChangeLog:

	* gimplify.cc (gimplify_scan_omp_clauses): Handle
	OpenMP device_type clause.
	* omp-low.cc (scan_sharing_clauses): Likewise.
	(lower_omp_target): Print 'sorry, unimplemented' for
	device_type clause value other than 'any'.

gcc/testsuite/ChangeLog:

	* c-c++-common/gomp/target-device-type-1.c: New test.
	* gfortran.dg/gomp/target-device-type-1.f90: New test.

diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc
index 9dd0d1450dd8383f15ded81f8cf55384e415e0cc..69b01253c0cd91ff7ce2261bb265a4e3cf410bfd 100644 (file)
--- a/gcc/c-family/c-omp.cc
+++ b/gcc/c-family/c-omp.cc
@@ -2176,6 +2176,7 @@ c_omp_split_clauses (location_t loc, enum tree_code code,
        {
        /* First the clauses that are unique to some constructs.  */
        case OMP_CLAUSE_DEVICE:
+       case OMP_CLAUSE_DEVICE_TYPE:
        case OMP_CLAUSE_DEFAULTMAP:
        case OMP_CLAUSE_DEPEND:
        case OMP_CLAUSE_DYN_GROUPPRIVATE:

diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 6e3c7bea3bc8cd10dcbaa7a6e0b93c7434f6b802..d6d0b0ed41517e39cc9fa2df66f18b7d241cc745 100644 (file)
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -27217,6 +27217,7 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser,
        | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEFAULTMAP)   \
        | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEPEND)       \
        | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE)       \
+       | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE_TYPE)  \
        | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DYN_GROUPPRIVATE) \
        | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_FIRSTPRIVATE) \
        | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR) \

diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index d9c8cfada9354bb8a327d97cfa28984747e40f43..18df3b1014e9a4c8ed0f3a709d18af6da005aa15 100644 (file)
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -50663,6 +50663,7 @@ cp_parser_omp_target_update (cp_parser *parser, cp_token *pragma_tok,
        | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEFAULTMAP)   \
        | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEPEND)       \
        | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE)       \
+       | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE_TYPE)  \
        | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DYN_GROUPPRIVATE) \
        | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_FIRSTPRIVATE) \
        | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR) \

diff --git a/gcc/fortran/dump-parse-tree.cc b/gcc/fortran/dump-parse-tree.cc
index 2a4ebb0fa0f91720cb8beb2e118390eba4dbd954..db6d54f5fc741d0fdf46af1ce678317683c8baa2 100644 (file)
--- a/gcc/fortran/dump-parse-tree.cc
+++ b/gcc/fortran/dump-parse-tree.cc
@@ -1984,6 +1984,20 @@ show_omp_clauses (gfc_omp_clauses *omp_clauses)
     fputs (" NOWAIT", dumpfile);
   if (omp_clauses->collapse)
     fprintf (dumpfile, " COLLAPSE(%d)", omp_clauses->collapse);
+  if (omp_clauses->device_type != OMP_DEVICE_TYPE_UNSET)
+    {
+      const char *s;
+      switch (omp_clauses->device_type)
+       {
+       case OMP_DEVICE_TYPE_HOST: s = "host"; break;
+       case OMP_DEVICE_TYPE_NOHOST: s = "nohost"; break;
+       case OMP_DEVICE_TYPE_ANY: s = "any"; break;
+       case OMP_DEVICE_TYPE_UNSET: gcc_unreachable ();
+       }
+      fputs (" DEVICE_TYPE(", dumpfile);
+      fputs (s, dumpfile);
+      fputc (')', dumpfile);
+    }
   for (list_type = 0; list_type < OMP_LIST_NUM; list_type++)
     if (omp_clauses->lists[list_type] != NULL)
       {

diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index af89c87b0ab0732e8019196c510a881375ce4528..4527068f97450eddc1bab295f33e59beecf223d6 100644 (file)
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -3013,6 +3013,43 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
                                           OMP_MAP_DETACH, false,
                                           allow_derived))
            continue;
+         if ((mask & OMP_CLAUSE_DEVICEPTR)
+             && gfc_match ("deviceptr ( ") == MATCH_YES
+             && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
+                                          OMP_MAP_FORCE_DEVICEPTR, false,
+                                          allow_derived))
+           continue;
+         if ((mask & OMP_CLAUSE_DEVICE_TYPE)
+             && gfc_match_dupl_check (c->device_type == OMP_DEVICE_TYPE_UNSET,
+                                      "device_type", true) == MATCH_YES)
+           {
+             if (gfc_match ("host") == MATCH_YES)
+               c->device_type = OMP_DEVICE_TYPE_HOST;
+             else if (gfc_match ("nohost") == MATCH_YES)
+               c->device_type = OMP_DEVICE_TYPE_NOHOST;
+             else if (gfc_match ("any") == MATCH_YES)
+               c->device_type = OMP_DEVICE_TYPE_ANY;
+             else
+               {
+                 gfc_error ("Expected HOST, NOHOST or ANY at %C");
+                 break;
+               }
+             if (gfc_match (" )") != MATCH_YES)
+               break;
+             continue;
+           }
+         if ((mask & OMP_CLAUSE_DEVICE_RESIDENT)
+             && gfc_match_omp_variable_list
+                  ("device_resident (",
+                   &c->lists[OMP_LIST_DEVICE_RESIDENT], true) == MATCH_YES)
+           continue;
+         if ((mask & OMP_CLAUSE_DEVICE)
+             && openacc
+             && gfc_match ("device ( ") == MATCH_YES
+             && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
+                                          OMP_MAP_FORCE_TO, true,
+                                          /* allow_derived = */ true))
+           continue;
          if ((mask & OMP_CLAUSE_DEVICE)
              && !openacc
              && ((m = gfc_match_dupl_check (!c->device, "device", true))
@@ -3072,42 +3109,6 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
                }
              continue;
            }
-         if ((mask & OMP_CLAUSE_DEVICE)
-             && openacc
-             && gfc_match ("device ( ") == MATCH_YES
-             && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
-                                          OMP_MAP_FORCE_TO, true,
-                                          /* allow_derived = */ true))
-           continue;
-         if ((mask & OMP_CLAUSE_DEVICEPTR)
-             && gfc_match ("deviceptr ( ") == MATCH_YES
-             && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
-                                          OMP_MAP_FORCE_DEVICEPTR, false,
-                                          allow_derived))
-           continue;
-         if ((mask & OMP_CLAUSE_DEVICE_TYPE)
-             && gfc_match ("device_type ( ") == MATCH_YES)
-           {
-             if (gfc_match ("host") == MATCH_YES)
-               c->device_type = OMP_DEVICE_TYPE_HOST;
-             else if (gfc_match ("nohost") == MATCH_YES)
-               c->device_type = OMP_DEVICE_TYPE_NOHOST;
-             else if (gfc_match ("any") == MATCH_YES)
-               c->device_type = OMP_DEVICE_TYPE_ANY;
-             else
-               {
-                 gfc_error ("Expected HOST, NOHOST or ANY at %C");
-                 break;
-               }
-             if (gfc_match (" )") != MATCH_YES)
-               break;
-             continue;
-           }
-         if ((mask & OMP_CLAUSE_DEVICE_RESIDENT)
-             && gfc_match_omp_variable_list
-                  ("device_resident (",
-                   &c->lists[OMP_LIST_DEVICE_RESIDENT], true) == MATCH_YES)
-           continue;
          if ((mask & OMP_CLAUSE_DIST_SCHEDULE)
              && c->dist_sched_kind == OMP_SCHED_NONE
              && gfc_match ("dist_schedule ( static") == MATCH_YES)
@@ -5136,7 +5137,7 @@ cleanup:
    | OMP_CLAUSE_IS_DEVICE_PTR | OMP_CLAUSE_IN_REDUCTION                        \
    | OMP_CLAUSE_THREAD_LIMIT | OMP_CLAUSE_ALLOCATE                     \
    | OMP_CLAUSE_HAS_DEVICE_ADDR | OMP_CLAUSE_USES_ALLOCATORS           \
-   | OMP_CLAUSE_DYN_GROUPPRIVATE)
+   | OMP_CLAUSE_DYN_GROUPPRIVATE | OMP_CLAUSE_DEVICE_TYPE)
 #define OMP_TARGET_DATA_CLAUSES \
   (omp_mask (OMP_CLAUSE_DEVICE) | OMP_CLAUSE_MAP | OMP_CLAUSE_IF       \
    | OMP_CLAUSE_USE_DEVICE_PTR | OMP_CLAUSE_USE_DEVICE_ADDR)

diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index a07dc2ec0e98f7c9ceb724b8d92f01c426885ccc..254fc934af1914dcf28262dd15af206f85b08b97 100644 (file)
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -5266,6 +5266,28 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
       omp_clauses = gfc_trans_add_clause (c, omp_clauses);
     }
 
+  if (clauses->device_type != OMP_DEVICE_TYPE_UNSET)
+    {
+      enum omp_clause_device_type_kind type;
+      switch (clauses->device_type)
+       {
+       case OMP_DEVICE_TYPE_HOST:
+         type = OMP_CLAUSE_DEVICE_TYPE_HOST;
+         break;
+       case OMP_DEVICE_TYPE_NOHOST:
+         type = OMP_CLAUSE_DEVICE_TYPE_NOHOST;
+         break;
+       case OMP_DEVICE_TYPE_ANY:
+         type = OMP_CLAUSE_DEVICE_TYPE_ANY;
+         break;
+       case OMP_DEVICE_TYPE_UNSET:
+         gcc_unreachable ();
+       }
+      c = build_omp_clause (gfc_get_location (&where), OMP_CLAUSE_DEVICE_TYPE);
+      OMP_CLAUSE_DEVICE_TYPE_KIND (c) = type;
+      omp_clauses = gfc_trans_add_clause (c, omp_clauses);
+    }
+
   if (clauses->dyn_groupprivate)
     {
       gfc_init_se (&se, NULL);
@@ -8051,6 +8073,8 @@ gfc_split_omp_clauses (gfc_code *code,
            = code->ext.omp_clauses->if_expr;
          clausesa[GFC_OMP_SPLIT_TARGET].nowait
            = code->ext.omp_clauses->nowait;
+         clausesa[GFC_OMP_SPLIT_TARGET].device_type
+           = code->ext.omp_clauses->device_type;
        }
       if (mask & GFC_OMP_MASK_TEAMS)
        {

diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index ee565336b1d72d18ea44c2c348e7f84d64d00852..751b4697271440ee29541d690eb677c14f3cc09e 100644 (file)
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -14885,6 +14885,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
        case OMP_CLAUSE_INIT:
        case OMP_CLAUSE_USE:
        case OMP_CLAUSE_DESTROY:
+       case OMP_CLAUSE_DEVICE_TYPE:
          break;
 
        case OMP_CLAUSE_DYN_GROUPPRIVATE:
@@ -16373,6 +16374,7 @@ end_adjust_omp_map_clause:
        case OMP_CLAUSE_INCLUSIVE:
        case OMP_CLAUSE_EXCLUSIVE:
        case OMP_CLAUSE_USES_ALLOCATORS:
+       case OMP_CLAUSE_DEVICE_TYPE:
          break;
 
        case OMP_CLAUSE_NOHOST:

diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 6fd685cdecd3ff8f7dc242bfd287ea5fbef0ee97..4540a25d1ad6f60240243be7dd545c3d874e8b48 100644 (file)
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -1768,6 +1768,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
        case OMP_CLAUSE_FINALIZE:
        case OMP_CLAUSE_TASK_REDUCTION:
        case OMP_CLAUSE_ALLOCATE:
+       case OMP_CLAUSE_DEVICE_TYPE:
          break;
 
        case OMP_CLAUSE_ALIGNED:
@@ -1994,6 +1995,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
        case OMP_CLAUSE_INIT:
        case OMP_CLAUSE_USE:
        case OMP_CLAUSE_DESTROY:
+       case OMP_CLAUSE_DEVICE_TYPE:
          break;
 
        case OMP_CLAUSE__CACHE_:
@@ -13098,6 +13100,13 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
            DECL_HAS_VALUE_EXPR_P (new_var) = 1;
          }
        break;
+       case OMP_CLAUSE_DEVICE_TYPE:
+         /* FIXME: Ensure that 'nohost' also has not implied before that
+            'g->have_offload = true' or an implicit declare target.  */
+         if (OMP_CLAUSE_DEVICE_TYPE_KIND (c) != OMP_CLAUSE_DEVICE_TYPE_ANY)
+           sorry_at (OMP_CLAUSE_LOCATION (c),
+                     "only the %<device_type(any)%> is supported");
+         break;
       }
 
   if (offloaded)

diff --git a/gcc/testsuite/c-c++-common/gomp/target-device-type-1.c b/gcc/testsuite/c-c++-common/gomp/target-device-type-1.c
new file mode 100644 (file)
index 0000000..e64349b
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-device-type-1.c
@@ -0,0 +1,24 @@
+// { dg-do compile }
+// { dg-additional-options "-fdump-tree-gimple" }
+
+void f ()
+{
+
+#pragma omp target
+  ;
+
+#pragma omp target device_type ( any )
+  ;
+
+#pragma omp target device_type ( nohost )  // { dg-message "sorry, unimplemented: only the 'device_type\\(any\\)' is supported" }
+  ;
+
+#pragma omp target device_type ( host )
+  ;
+
+}
+
+// { dg-final { scan-tree-dump-times "#pragma omp target num_teams\\(-2\\) thread_limit\\(0\\)\[\\r\\n\]" 1 "gimple" } }
+// { dg-final { scan-tree-dump-times "#pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) device_type\\(any\\)" 1 "gimple" } }
+// { dg-final { scan-tree-dump-times "#pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) device_type\\(nohost\\)" 1 "gimple" } }
+// { dg-final { scan-tree-dump-times "#pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) device_type\\(host\\)" 1 "gimple" } }

diff --git a/gcc/testsuite/gfortran.dg/gomp/target-device-type-1.f90 b/gcc/testsuite/gfortran.dg/gomp/target-device-type-1.f90
new file mode 100644 (file)
index 0000000..be33bb6
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/target-device-type-1.f90
@@ -0,0 +1,21 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+!$omp target
+!$omp end target
+
+!$omp target device_type ( any )
+!$omp end target
+
+!$omp target device_type ( nohost )  ! { dg-message "sorry, unimplemented: only the 'device_type\\(any\\)' is supported" }
+!$omp end target
+
+!$omp target device_type ( host )
+!$omp end target
+
+end
+
+! { dg-final { scan-tree-dump-times "#pragma omp target num_teams\\(-2\\) thread_limit\\(0\\)\[\\r\\n\]" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) device_type\\(any\\)" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) device_type\\(nohost\\)" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) device_type\\(host\\)" 1 "gimple" } }