diff mbox series

OpenMP: 'interop' construct - add C parser support, improve Fortran pasing

Message ID b5725d74-8ba2-46d1-98c8-774cc202b52f@baylibre.com
State New
Headers show
Series OpenMP: 'interop' construct - add C parser support, improve Fortran pasing | expand

Commit Message

Tobias Burnus Nov. 11, 2024, 3:01 p.m. UTC
Background:
   omp interop device(1) init(prefer_type("cuda"), targetsync: obj) depend(inout: x) nowait
     …
   omp interop destroy(obj)

initializes the omp_interop_t / integer(omp_interop_kind) variable for device '1'
and (thanks to 'targetsync') creates a stream object. 'obj' can then be used as
follows: first, we have to check the returned type (e.g. CUDA as wished for or
something else including N/A alias omp_interop_fr_none).
The (CUDA) stream, (cuda) device number etc. can then be extracted and used with
CUDA calls.

In terms of the parser, that's quite boring if there wasn't the prefer_type modifier.
Besides taking a list of strings and constant integer expressions, OpenMP 6.0 also
permits:
prefer_type( {fr("cuda"), attr("ompx_1", "ompx_2")},
              {attr("ompx_2"), attr("ompx_4")} )
i.e. the same to 'fr' and a string to 'attr' that must start with 'ompx_'. There
can be 0 or 1 'fr' and >= 0 'attr' per curly brace (but at least one 'fr'/'attr'
must be specified).

* * *

The attached patch add the C parser, which in turn means that there is now a
middle-end representation for it. Additionally, it fixes the Fortran compiler
for issues found while doing so - and for a newer OpenMP 6 spec change, i.e.
only one 'fr' value permitted per {…} and the constant integer value to 'fr'
may be any const integer expr not only an identifier, which is the same as
for the old, simpler 'prefer_type("hip",int_expr,"sycl")' syntax.

Comments, remarks, concerns, suggestions before I commit it?

Tobias

* * *

PS: Once 'omp interop' has returned the object, the
https://gcc.gnu.org/onlinedocs/libgomp/Interoperability-Routines.html
can be used to access it. See libgomp.*/interop-routines*.{F,F90,c} for
some testcases. - Proper combined testcases will be added once the
compiler middle-end + libgomp parts have been implemented to connect the
two. → TODO: C++ parser, middle-end code including calling new libgomp function.

Once done, the AMD GPU (gcn) and Nvidia GPU libgomp plugins need to handle
it to return an interop object for CUDA/CUDA_DRIVER/HIP/HSA; I posted an
RFC patch the other day, which should mostly work once (↑) is done; it still
requires some updates, cleanups and additions, but otherwise … :-)
https://gcc.gnu.org/pipermail/gcc-patches/2024-August/661207.html
diff mbox series

Patch

OpenMP: 'interop' construct - add C parser support, improve Fortran pasing

Add middle end support for the 'interop' directive and the 'init', 'use',
and 'destroy' clauses - but fail with a sorry, unimplemented in gimplify.cc.

For Fortran, generate the tree code, update the internal representation,
add some more diagnostic checks and update for newer specification changes
('fr' only takes a single value, but it integer expressions are permitted
again [like with the old syntax] not only constant identifiers).

For C, this patch adds the full parser support for 'interop'.

Still missing (later commit) is parsing support in C++ and actually
handling the directive in the middle end and in libgomp.

The GOMP_INTEROP_IFR_* internal values have been changed to have space
for vendor specific values that are adjacent to the existing values
but negative, if needed.

gcc/c-family/ChangeLog:

	* c-common.h (c_omp_interop_t_p): New prototype.
	* c-omp.cc (c_omp_interop_t_p): Check whether the type is
	omp_interop_t.
	(c_omp_directives): Uncomment 'interop'.
	* c-pragma.cc (omp_pragmas): Add 'interop'.
	* c-pragma.h (enum pragma_kind): Add PRAGMA_OMP_INTEROP.
	(enum pragma_omp_clause): Add init, use, and destroy clauses.

gcc/c/ChangeLog:

	* c-parser.cc (INCLUDE_STRING): Define.
	(c_parser_pragma): Handle 'interop' directive.
	(c_parser_omp_clause_name, c_parser_omp_all_clauses): Handle init,
	use, and destroy clauses.
	(c_parser_omp_clause_destroy, c_parser_omp_modifier_prefer_type,
	c_parser_omp_clause_init, c_parser_omp_clause_use,
	OMP_INTEROP_CLAUSE_MASK, c_parser_omp_interop): New.
	* c-typeck.cc (c_finish_omp_clauses): Add missing OPT_Wopenmp to
	a warning; handle new clauses.

gcc/fortran/ChangeLog:

	* gfortran.h (gfc_omp_namelist): Cleanup interop internal
	representation.
	* dump-parse-tree.cc (show_omp_namelist): Update for changed
	internal representation.
	* match.cc (gfc_free_omp_namelist): Likewise.
	* openmp.cc (gfc_match_omp_prefer_type, gfc_match_omp_init):
	Likewise; also handle some corner cases better and update for
	newer 6.0 changes related to 'fr'.
	(resolve_omp_clauses): Add type-check for interop variables.
	* trans-openmp.cc (gfc_trans_omp_clauses): Handle init, use
	and destroy clauses.
	(gfc_trans_openmp_interop): New.
	(gfc_trans_omp_directive): Call it.

gcc/ChangeLog:

	* gimplify.cc (gimplify_expr): Handle OMP_INTEROP by printing
	"sorry, uninplemented".
	* omp-general.cc (omp_get_fr_id_from_name): Return
	GOMP_INTEROP_IFR_UNKNOWN not 0 if not found.
	(omp_get_name_from_fr_id): Return "<unknown>" not NULL
	if not found (used for dumps).
	* tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_DESTROY,
	OMP_CLAUSE_USE, and OMP_CLAUSE_INIT.
	* tree-pretty-print.cc (dump_omp_init_prefer_type): New.
	(dump_omp_clause): Handle init, use and destroy clauses.
	(dump_generic_node): Handle interop directive.
	* tree.cc (omp_clause_num_ops, omp_clause_code_name): Add new
	init/use/destroy clauses.
	* tree.def (OACC_LOOP): Fix comment.
	(OMP_INTEROP): Add.
	* tree.h (OMP_INTEROP_CLAUSES, OMP_CLAUSE_INIT_TARGET,
	OMP_CLAUSE_INIT_TARGETSYNC, OMP_CLAUSE_INIT_PREFER_TYPE): New.

include/ChangeLog:

	* gomp-constants.h (GOMP_INTEROP_IFR_NONE): Rename ...
	(GOMP_INTEROP_IFR_UNKNOWN): ... to this. And change value.
	(GOMP_INTEROP_IFR_SEPARATOR): Likewise.

gcc/testsuite/ChangeLog:

	* gfortran.dg/gomp/interop-1.f90: Update for parser changes,
	spec changes and add new tests.
	* gfortran.dg/gomp/interop-2.f90: Likewise.
	* gfortran.dg/gomp/interop-3.f90: Likewise.
	* c-c++-common/gomp/interop-1.c: New test.
	* c-c++-common/gomp/interop-2.c: New test.
	* c-c++-common/gomp/interop-3.c: New test.
	* c-c++-common/gomp/interop-4.c: New test.
	* gfortran.dg/gomp/interop-4.f90: New test.

 gcc/c-family/c-common.h                      |   1 +
 gcc/c-family/c-omp.cc                        |  22 +-
 gcc/c-family/c-pragma.cc                     |   1 +
 gcc/c-family/c-pragma.h                      |   4 +
 gcc/c/c-parser.cc                            | 478 +++++++++++++++++++++++++++
 gcc/c/c-typeck.cc                            |  15 +-
 gcc/fortran/dump-parse-tree.cc               |  69 ++--
 gcc/fortran/gfortran.h                       |   3 +-
 gcc/fortran/match.cc                         |   9 +-
 gcc/fortran/openmp.cc                        | 234 ++++++-------
 gcc/fortran/trans-openmp.cc                  |  59 +++-
 gcc/gimplify.cc                              |   5 +
 gcc/omp-general.cc                           |   4 +-
 gcc/testsuite/c-c++-common/gomp/interop-1.c  | 115 +++++++
 gcc/testsuite/c-c++-common/gomp/interop-2.c  | 106 ++++++
 gcc/testsuite/c-c++-common/gomp/interop-3.c  |  81 +++++
 gcc/testsuite/c-c++-common/gomp/interop-4.c  |  76 +++++
 gcc/testsuite/gfortran.dg/gomp/interop-1.f90 |  41 ++-
 gcc/testsuite/gfortran.dg/gomp/interop-2.f90 |  40 ++-
 gcc/testsuite/gfortran.dg/gomp/interop-3.f90 |  19 +-
 gcc/testsuite/gfortran.dg/gomp/interop-4.f90 |  56 ++++
 gcc/tree-core.h                              |  13 +
 gcc/tree-pretty-print.cc                     |  87 +++++
 gcc/tree.cc                                  |   6 +
 gcc/tree.def                                 |   6 +-
 gcc/tree.h                                   |  12 +
 include/gomp-constants.h                     |   7 +-
 27 files changed, 1367 insertions(+), 202 deletions(-)

diff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h
index 46099b63514..25466c61bc2 100644
--- a/gcc/c-family/c-common.h
+++ b/gcc/c-family/c-common.h
@@ -1310,6 +1310,7 @@  extern void c_finish_omp_barrier (location_t);
 extern tree c_finish_omp_atomic (location_t, enum tree_code, enum tree_code,
 				 tree, tree, tree, tree, tree, tree, bool,
 				 enum omp_memory_order, bool, bool = false);
+extern bool c_omp_interop_t_p (tree);
 extern bool c_omp_depend_t_p (tree);
 extern void c_finish_omp_depobj (location_t, tree, enum omp_clause_depend_kind,
 				 tree);
diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc
index eacfcdc3437..6062bd9665d 100644
--- a/gcc/c-family/c-omp.cc
+++ b/gcc/c-family/c-omp.cc
@@ -664,6 +664,24 @@  c_finish_omp_atomic (location_t loc, enum tree_code code,
 }
 
 
+/* Return true if TYPE is the implementation's omp_interop_t.  */
+
+bool
+c_omp_interop_t_p (tree type)
+{
+  type = TYPE_MAIN_VARIANT (type);
+  return (TREE_CODE (type) == ENUMERAL_TYPE
+	  && TYPE_NAME (type)
+	  && ((TREE_CODE (TYPE_NAME (type)) == TYPE_DECL
+	       ? DECL_NAME (TYPE_NAME (type)) : TYPE_NAME (type))
+	      == get_identifier ("omp_interop_t"))
+	  && TYPE_FILE_SCOPE_P (type)
+	  && COMPLETE_TYPE_P (type)
+	  && TREE_CODE (TYPE_SIZE (type)) == INTEGER_CST
+	  && !compare_tree_int (TYPE_SIZE (type),
+				tree_to_uhwi (TYPE_SIZE (ptr_type_node))));
+}
+
 /* Return true if TYPE is the implementation's omp_depend_t.  */
 
 bool
@@ -4321,8 +4339,8 @@  const struct c_omp_directive c_omp_directives[] = {
     C_OMP_DIR_CONSTRUCT, true },
   /* { "groupprivate", nullptr, nullptr, PRAGMA_OMP_GROUPPRIVATE,
     C_OMP_DIR_DECLARATIVE, false },  */
-  /* { "interop", nullptr, nullptr, PRAGMA_OMP_INTEROP,
-    C_OMP_DIR_STANDALONE, false },  */
+  { "interop", nullptr, nullptr, PRAGMA_OMP_INTEROP,
+    C_OMP_DIR_STANDALONE, false },
   { "loop", nullptr, nullptr, PRAGMA_OMP_LOOP,
     C_OMP_DIR_CONSTRUCT, true },
   { "masked", nullptr, nullptr, PRAGMA_OMP_MASKED,
diff --git a/gcc/c-family/c-pragma.cc b/gcc/c-family/c-pragma.cc
index c78721824e3..ea7446a4abe 100644
--- a/gcc/c-family/c-pragma.cc
+++ b/gcc/c-family/c-pragma.cc
@@ -1528,6 +1528,7 @@  static const struct omp_pragma_def omp_pragmas[] = {
   { "error", PRAGMA_OMP_ERROR },
   { "end", PRAGMA_OMP_END },
   { "flush", PRAGMA_OMP_FLUSH },
+  { "interop", PRAGMA_OMP_INTEROP },
   { "nothing", PRAGMA_OMP_NOTHING },
   { "requires", PRAGMA_OMP_REQUIRES },
   { "scope", PRAGMA_OMP_SCOPE },
diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index 2ebde06c471..715d3f077a2 100644
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -60,6 +60,7 @@  enum pragma_kind {
   PRAGMA_OMP_END,
   PRAGMA_OMP_FLUSH,
   PRAGMA_OMP_FOR,
+  PRAGMA_OMP_INTEROP,
   PRAGMA_OMP_LOOP,
   PRAGMA_OMP_NOTHING,
   PRAGMA_OMP_MASKED,
@@ -110,6 +111,7 @@  enum pragma_omp_clause {
   PRAGMA_OMP_CLAUSE_DEFAULT,
   PRAGMA_OMP_CLAUSE_DEFAULTMAP,
   PRAGMA_OMP_CLAUSE_DEPEND,
+  PRAGMA_OMP_CLAUSE_DESTROY,
   PRAGMA_OMP_CLAUSE_DETACH,
   PRAGMA_OMP_CLAUSE_DEVICE,
   PRAGMA_OMP_CLAUSE_DEVICE_TYPE,
@@ -129,6 +131,7 @@  enum pragma_omp_clause {
   PRAGMA_OMP_CLAUSE_IN_REDUCTION,
   PRAGMA_OMP_CLAUSE_INBRANCH,
   PRAGMA_OMP_CLAUSE_INDIRECT,
+  PRAGMA_OMP_CLAUSE_INIT,
   PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR,
   PRAGMA_OMP_CLAUSE_LASTPRIVATE,
   PRAGMA_OMP_CLAUSE_LINEAR,
@@ -163,6 +166,7 @@  enum pragma_omp_clause {
   PRAGMA_OMP_CLAUSE_TO,
   PRAGMA_OMP_CLAUSE_UNIFORM,
   PRAGMA_OMP_CLAUSE_UNTIED,
+  PRAGMA_OMP_CLAUSE_USE,
   PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR,
   PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR,
 
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 3ab8a49bf35..b08ca05d143 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -37,6 +37,7 @@  along with GCC; see the file COPYING3.  If not see
 
 #include "config.h"
 #define INCLUDE_MEMORY
+#define INCLUDE_STRING
 #include "system.h"
 #include "coretypes.h"
 #include "target.h"
@@ -1750,6 +1751,7 @@  static void c_parser_omp_assumption_clauses (c_parser *, bool);
 static void c_parser_omp_allocate (c_parser *);
 static void c_parser_omp_assumes (c_parser *);
 static bool c_parser_omp_ordered (c_parser *, enum pragma_context, bool *);
+static void c_parser_omp_interop (c_parser *);
 static void c_parser_oacc_routine (c_parser *, enum pragma_context);
 
 /* These Objective-C parser functions are only ever called when
@@ -15343,6 +15345,15 @@  c_parser_pragma (c_parser *parser, enum pragma_context context, bool *if_p,
       c_parser_omp_flush (parser);
       return false;
 
+    case PRAGMA_OMP_INTEROP:
+      if (context != pragma_compound)
+	{
+	  construct = "omp interop";
+	  goto in_compound;
+	}
+      c_parser_omp_interop (parser);
+      return false;
+
     case PRAGMA_OMP_TASKWAIT:
       if (context != pragma_compound)
 	{
@@ -15670,6 +15681,8 @@  c_parser_omp_clause_name (c_parser *parser)
 	    result = PRAGMA_OACC_CLAUSE_DELETE;
 	  else if (!strcmp ("depend", p))
 	    result = PRAGMA_OMP_CLAUSE_DEPEND;
+	  else if (!strcmp ("destroy", p))
+	    result = PRAGMA_OMP_CLAUSE_DESTROY;
 	  else if (!strcmp ("detach", p))
 	    result = PRAGMA_OACC_CLAUSE_DETACH;
 	  else if (!strcmp ("device", p))
@@ -15728,6 +15741,8 @@  c_parser_omp_clause_name (c_parser *parser)
 	    result = PRAGMA_OACC_CLAUSE_INDEPENDENT;
 	  else if (!strcmp ("indirect", p))
 	    result = PRAGMA_OMP_CLAUSE_INDIRECT;
+	  else if (!strcmp ("init", p))
+	    result = PRAGMA_OMP_CLAUSE_INIT;
 	  else if (!strcmp ("is_device_ptr", p))
 	    result = PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR;
 	  break;
@@ -15844,6 +15859,8 @@  c_parser_omp_clause_name (c_parser *parser)
 	    result = PRAGMA_OMP_CLAUSE_UNIFORM;
 	  else if (!strcmp ("untied", p))
 	    result = PRAGMA_OMP_CLAUSE_UNTIED;
+	  else if (!strcmp ("use", p))
+	    result = PRAGMA_OMP_CLAUSE_USE;
 	  else if (!strcmp ("use_device", p))
 	    result = PRAGMA_OACC_CLAUSE_USE_DEVICE;
 	  else if (!strcmp ("use_device_addr", p))
@@ -20073,6 +20090,386 @@  c_parser_omp_clause_detach (c_parser *parser, tree list)
   return u;
 }
 
+/* OpenMP 5.0:
+   destroy ( variable-list ) */
+
+static tree
+c_parser_omp_clause_destroy (c_parser *parser, tree list)
+{
+  return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_DESTROY, list);
+}
+
+/* OpenMP 5.1:
+     prefer_type ( const-int-expr-or-string-literal-list )
+
+   OpenMP 6.0:
+     prefer_type ( { preference-selector-list }, { ... } )
+
+   with preference-selector being:
+     fr ( identifier-or-string-literal-list )
+     attr ( string-list )
+
+   Data format:
+    For the foreign runtime identifiers, string values are converted to
+    their integer value; unknown string or integer values are set to
+    GOMP_INTEROP_IFR_KNOWN.
+
+    Each item (a) GOMP_INTEROP_IFR_SEPARATOR
+	      (b) for any 'fr', its integer value.
+		  Note: Spec only permits 1 'fr' entry (6.0; changed after TR13)
+	      (c) GOMP_INTEROP_IFR_SEPARATOR
+	      (d) list of \0-terminated non-empty strings for 'attr'
+	      (e) '\0'
+    Tailing '\0'.  */
+
+static tree
+c_parser_omp_modifier_prefer_type (c_parser *parser)
+{
+  if (!c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>"))
+    return error_mark_node;
+
+  std::string str;
+
+  /* Old Format: const-int-expr-or-string-literal-list */
+  if (!c_parser_next_token_is (parser, CPP_OPEN_BRACE))
+    while (true)
+      {
+	str += (char) GOMP_INTEROP_IFR_SEPARATOR;
+	if (c_parser_next_token_is (parser, CPP_STRING))
+	  {
+	    c_expr cval = c_parser_string_literal (parser, false, false);
+	    if (cval.value == error_mark_node)
+	      return error_mark_node;
+	    if ((size_t) TREE_STRING_LENGTH (cval.value)
+		!= strlen (TREE_STRING_POINTER (cval.value)) + 1)
+	      {
+		error_at (cval.get_location (), "string literal must "
+						"not contain %<\\0%>");
+		parser->error = true;
+		return error_mark_node;
+	      }
+
+	    char c = omp_get_fr_id_from_name (TREE_STRING_POINTER (cval.value));
+	    if (c == GOMP_INTEROP_IFR_UNKNOWN)
+	      warning_at (cval.get_location (), OPT_Wopenmp,
+			  "unknown foreign runtime identifier %qs",
+			  TREE_STRING_POINTER (cval.value));
+	    str += c;
+	  }
+	else
+	  {
+	    c_expr cval = c_parser_expr_no_commas (parser, NULL);
+	    tree value = c_fully_fold (cval.value, false, NULL);
+	    if (INTEGRAL_TYPE_P (TREE_TYPE (value))
+		&& TREE_CODE (value) != INTEGER_CST)
+	      value = convert_lvalue_to_rvalue (cval.get_start (), cval,
+						false, true).value;
+	    if (TREE_CODE (value) != INTEGER_CST
+		|| !tree_fits_shwi_p (value))
+	      {
+		c_parser_error (parser, "expected string literal or constant "
+					"integer expression");
+		return error_mark_node;
+	      }
+	    HOST_WIDE_INT n = tree_to_shwi (value);
+	    if (n < 1 || n > GOMP_INTEROP_IFR_LAST)
+	      {
+		warning_at (cval.get_location (), OPT_Wopenmp,
+			    "unknown foreign runtime identifier %qwd", n);
+		n = GOMP_INTEROP_IFR_UNKNOWN;
+	      }
+	    str += (char) n;
+	  }
+	str += (char) GOMP_INTEROP_IFR_SEPARATOR;
+	str += '\0';
+	if (c_parser_next_token_is (parser, CPP_COMMA))
+	  {
+	    c_parser_consume_token (parser);
+	    continue;
+	  }
+	if (!c_parser_require (parser, CPP_CLOSE_PAREN,
+			       "expected %<,%> or %<)%>"))
+	  return error_mark_node;
+	str += '\0';
+	tree res = build_string (str.length (), str.data ());
+	TREE_TYPE (res) = build_array_type_nelts (unsigned_char_type_node,
+						  str.length ());
+	return res;
+      }
+
+  /* New format. */
+  std::string str2;
+  while (true)
+    {
+      if (!c_parser_require (parser, CPP_OPEN_BRACE, "expected %<{%>"))
+	return error_mark_node;
+      str += (char) GOMP_INTEROP_IFR_SEPARATOR;
+      str2.clear ();
+      bool has_fr = false;
+      while (true)
+	{
+	  c_token *tok = c_parser_peek_token (parser);
+	  if (tok->type != CPP_NAME
+	      || (strcmp("fr", IDENTIFIER_POINTER (tok->value)) != 0
+		  && strcmp("attr", IDENTIFIER_POINTER (tok->value)) != 0))
+	    {
+	      c_parser_error (parser, "expected %<fr%> or %<attr%> preference "
+				      "selector");
+	      return error_mark_node;
+	    }
+	  c_parser_consume_token (parser);
+	  bool is_fr = IDENTIFIER_POINTER (tok->value)[0] == 'f';
+	  if (is_fr && has_fr)
+	    {
+	      c_parser_error (parser, "duplicated %<fr%> preference selector");
+	      return error_mark_node;
+	    }
+	  if (!c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>"))
+	    return error_mark_node;
+	  while (true)
+	    {
+	      if (c_parser_next_token_is (parser, CPP_STRING))
+		{
+		  c_expr cval = c_parser_string_literal (parser, false, false);
+		  tree value = cval.value;
+		  if (value == error_mark_node)
+		    return error_mark_node;
+		  if ((size_t) TREE_STRING_LENGTH (value)
+		      != strlen (TREE_STRING_POINTER (value)) + 1)
+		    {
+		      error_at (cval.get_location (), "string literal must "
+						      "not contain %<\\0%>");
+		      parser->error = true;
+		      return error_mark_node;
+		    }
+		  if (!is_fr)
+		    {
+		      if (!startswith (TREE_STRING_POINTER (value), "ompx_"))
+			{
+			  error_at (cval.get_location (),
+				    "%<attr%> string literal must start with "
+				    "%<ompx_%>");
+			  parser->error = true;
+			  return error_mark_node;
+			}
+		      if (strchr (TREE_STRING_POINTER (value), ','))
+			{
+			  error_at (cval.get_location (),
+				    "%<attr%> string literal must not contain "
+				    "a comma");
+			  parser->error = true;
+			  return error_mark_node;
+			}
+		      str2 += TREE_STRING_POINTER (value);
+		      str2 += '\0';
+		    }
+		  else
+		    {
+		      if (*TREE_STRING_POINTER (value) == '\0')
+			{
+			  c_parser_error (parser, "non-empty string literal expected");
+			  return error_mark_node;
+			}
+		      char c = omp_get_fr_id_from_name (TREE_STRING_POINTER (value));
+		      if (c == GOMP_INTEROP_IFR_UNKNOWN)
+			warning_at (cval.get_location (), OPT_Wopenmp,
+				    "unknown foreign runtime identifier %qs",
+				    TREE_STRING_POINTER (value));
+		     str += c;
+		     has_fr = true;
+		    }
+		}
+	      else if (!is_fr)
+		{
+		  c_parser_error (parser, "expected string literal");
+		  return error_mark_node;
+		}
+	      else
+		{
+		  c_expr cval = c_parser_expr_no_commas (parser, NULL);
+		  tree value = c_fully_fold (cval.value, false, NULL);
+		  if (INTEGRAL_TYPE_P (TREE_TYPE (value))
+		      && TREE_CODE (value) != INTEGER_CST)
+		    value = convert_lvalue_to_rvalue (cval.get_start (), cval,
+						      false, true).value;
+
+		  if (TREE_CODE (value) != INTEGER_CST
+		      || !tree_fits_shwi_p (value))
+		    {
+		      c_parser_error (parser, "expected string literal or "
+					      "constant integer expression");
+		      return error_mark_node;
+		    }
+		  HOST_WIDE_INT n = tree_to_shwi (value);
+		  if (n < 1 || n > GOMP_INTEROP_IFR_LAST)
+		    {
+		      warning_at (cval.get_location (), OPT_Wopenmp,
+				  "unknown foreign runtime identifier %qwd", n);
+		      n = GOMP_INTEROP_IFR_UNKNOWN;
+		    }
+		  str += (char) n;
+		  has_fr = true;
+		}
+	      if (!is_fr
+		  && c_parser_next_token_is (parser, CPP_COMMA))
+		{
+		  c_parser_consume_token (parser);
+		  continue;
+		}
+	      if (!c_parser_require (parser, CPP_CLOSE_PAREN,
+				     is_fr ? G_("expected %<)%>")
+					   : G_("expected %<)%> or %<,%>")))
+		return error_mark_node;
+	      break;
+	    }
+	  if (c_parser_next_token_is (parser, CPP_COMMA))
+	    {
+	      c_parser_consume_token (parser);
+	      continue;
+	    }
+	  if (c_parser_next_token_is (parser, CPP_CLOSE_BRACE))
+	    break;
+	  c_parser_error (parser, "expected %<,%> or %<}%>");
+	  return error_mark_node;
+	}
+      str += (char) GOMP_INTEROP_IFR_SEPARATOR;
+      str += str2;
+      str += '\0';
+      c_parser_consume_token (parser);
+      if (c_parser_next_token_is (parser, CPP_CLOSE_PAREN))
+	break;
+      if (!c_parser_require (parser, CPP_COMMA, "expected %<)%> or %<,%>"))
+	return error_mark_node;
+    }
+  c_parser_consume_token (parser);
+  str += '\0';
+  tree res = build_string (str.length (), str.data ());
+  TREE_TYPE (res) = build_array_type_nelts (unsigned_char_type_node,
+						  str.length ());
+  return res;
+}
+
+/* OpenMP 5.1:
+   init ( [init-modifier-list : ] variable-list )
+
+   Modifiers:
+     target
+     targetsync
+     prefer_type (preference-specification) */
+
+static tree
+c_parser_omp_clause_init (c_parser *parser, tree list)
+{
+  location_t loc = c_parser_peek_token (parser)->location;
+
+  matching_parens parens;
+  if (!parens.require_open (parser))
+    return list;
+
+  unsigned pos = 0, raw_pos = 1;
+  while (c_parser_peek_nth_token_raw (parser, raw_pos)->type == CPP_NAME)
+    {
+      pos++; raw_pos++;
+      if (c_parser_peek_nth_token_raw (parser, raw_pos)->type == CPP_OPEN_PAREN)
+	{
+	  raw_pos++;
+	  c_parser_check_balanced_raw_token_sequence (parser, &raw_pos);
+	  if (c_parser_peek_nth_token_raw (parser, raw_pos)->type != CPP_CLOSE_PAREN)
+	    {
+	      pos = 0;
+	      break;
+	    }
+	  raw_pos++;
+	}
+      if (c_parser_peek_nth_token_raw (parser, raw_pos)->type == CPP_COLON)
+	break;
+      if (c_parser_peek_nth_token_raw (parser, raw_pos)->type != CPP_COMMA)
+	{
+	  pos = 0;
+	  break;
+	}
+      pos++;
+      raw_pos++;
+    }
+
+  bool target = false;
+  bool targetsync = false;
+  tree prefer_type_tree = NULL_TREE;
+
+  for (unsigned pos2 = 0; pos2 < pos; ++pos2)
+    {
+      c_token *tok = c_parser_peek_token (parser);
+      if (tok->type == CPP_COMMA)
+	{
+	  c_parser_consume_token (parser);
+	  continue;
+	}
+
+      const char *p = IDENTIFIER_POINTER (tok->value);
+      if (strcmp ("targetsync", p) == 0)
+	{
+	  if (targetsync)
+	    error_at (tok->location, "duplicate %<targetsync%> modifier");
+	  targetsync = true;
+	  c_parser_consume_token (parser);
+	}
+      else if (strcmp ("target", p) == 0)
+	{
+	  if (target)
+	    error_at (tok->location, "duplicate %<target%> modifier");
+	  target = true;
+	  c_parser_consume_token (parser);
+	}
+      else if (strcmp ("prefer_type", p) == 0)
+	{
+	  if (prefer_type_tree != NULL_TREE)
+	    error_at (tok->location, "duplicate %<prefer_type%> modifier");
+	  c_parser_consume_token (parser);
+	  prefer_type_tree = c_parser_omp_modifier_prefer_type (parser);
+	  if (prefer_type_tree == error_mark_node)
+	    return list;
+	}
+      else
+	{
+	  c_parser_error (parser, "%<init%> clause with modifier other than "
+				  "%<prefer_type%>, %<target%> or "
+				  "%<targetsync%>");
+	  parens.skip_until_found_close (parser);
+	  return list;
+	}
+    }
+  if (pos)
+    {
+      c_token *tok = c_parser_peek_token (parser);
+      gcc_checking_assert (tok->type == CPP_COLON);
+      c_parser_consume_token (parser);
+    }
+
+  tree nl = c_parser_omp_variable_list (parser, loc, OMP_CLAUSE_INIT, list,
+					 false);
+  parens.skip_until_found_close (parser);
+
+  for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
+    {
+      if (target)
+	OMP_CLAUSE_INIT_TARGET (c) = 1;
+      if (targetsync)
+	OMP_CLAUSE_INIT_TARGETSYNC (c) = 1;
+      if (prefer_type_tree)
+	OMP_CLAUSE_INIT_PREFER_TYPE (c) = prefer_type_tree;
+    }
+  return nl;
+}
+
+/* OpenMP 5.0:
+   use ( variable-list ) */
+
+static tree
+c_parser_omp_clause_use (c_parser *parser, tree list)
+{
+  return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_USE, list);
+}
+
 /* Parse all OpenACC clauses.  The set clauses allowed by the directive
    is a bitmask in MASK.  Return the list of clauses found.  */
 
@@ -20568,6 +20965,18 @@  c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask,
 	  clauses = c_parser_omp_clause_doacross (parser, clauses);
 	  c_name = "doacross";
 	  break;
+	case PRAGMA_OMP_CLAUSE_DESTROY:
+	  clauses = c_parser_omp_clause_destroy (parser, clauses);
+	  c_name = "destroy";
+	  break;
+	case PRAGMA_OMP_CLAUSE_INIT:
+	  clauses = c_parser_omp_clause_init (parser, clauses);
+	  c_name = "init";
+	  break;
+	case PRAGMA_OMP_CLAUSE_USE:
+	  clauses = c_parser_omp_clause_use (parser, clauses);
+	  c_name = "use";
+	  break;
 	case PRAGMA_OMP_CLAUSE_MAP:
 	  clauses = c_parser_omp_clause_map (parser, clauses);
 	  c_name = "map";
@@ -24002,6 +24411,75 @@  c_parser_omp_masked (location_t loc, c_parser *parser,
 			      clauses);
 }
 
+/* OpenMP 5.1:
+   # pragma omp interop clauses[opt] new-line */
+
+#define OMP_INTEROP_CLAUSE_MASK					\
+	( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEPEND)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DESTROY)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_INIT)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOWAIT)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USE))
+
+static void
+c_parser_omp_interop (c_parser *parser)
+{
+  location_t loc = c_parser_peek_token (parser)->location;
+  c_parser_consume_pragma (parser);
+  tree clauses = c_parser_omp_all_clauses (parser,
+					   OMP_INTEROP_CLAUSE_MASK,
+					   "#pragma omp interop");
+  tree depend_clause = NULL_TREE;
+  bool init_use_destroy_seen = false;
+  bool init_seen = true;
+  tree init_no_targetsync_clause = NULL_TREE;
+
+  for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+    switch (OMP_CLAUSE_CODE (c))
+      {
+	case OMP_CLAUSE_DEPEND:
+	  depend_clause = c;
+	  break;
+	case OMP_CLAUSE_INIT:
+	  init_seen = true;
+	  if (!OMP_CLAUSE_INIT_TARGETSYNC (c))
+	    init_no_targetsync_clause = c;
+	  /* FALLTHRU */
+	case OMP_CLAUSE_DESTROY:
+	case OMP_CLAUSE_USE:
+	  init_use_destroy_seen = true;
+	  if (!c_omp_interop_t_p (TREE_TYPE (OMP_CLAUSE_DECL (c))))
+	    error_at (OMP_CLAUSE_LOCATION (c),
+		      "%qD must be of %<omp_interop_t%>", OMP_CLAUSE_DECL (c));
+	  else if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_USE
+		   && TREE_READONLY (OMP_CLAUSE_DECL (c)))
+	    error_at (OMP_CLAUSE_LOCATION (c),
+		      "%qD shall not be const", OMP_CLAUSE_DECL (c));
+	  break;
+	default:
+	  break;
+      }
+
+  if (depend_clause
+      && (!init_use_destroy_seen
+	  || (init_seen && init_no_targetsync_clause)))
+    {
+      error_at (OMP_CLAUSE_LOCATION (depend_clause),
+		"%<depend%> clause requires action clauses with "
+		"%<targetsync%> interop-type");
+      if (init_no_targetsync_clause)
+	inform (OMP_CLAUSE_LOCATION (init_no_targetsync_clause),
+		"%<init%> clause lacks the %<targetsync%> modifier");
+    }
+
+  tree stmt = make_node (OMP_INTEROP);
+  TREE_TYPE (stmt) = void_type_node;
+  OMP_INTEROP_CLAUSES (stmt) = clauses;
+  SET_EXPR_LOCATION (stmt, loc);
+  add_stmt (stmt);
+}
+
 /* OpenMP 2.5:
    # pragma omp ordered new-line
      structured-block
diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc
index 201d75d2e9c..7a700ed1a49 100644
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -16219,7 +16219,7 @@  c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	    }
 	  else if (bitmap_bit_p (&aligned_head, DECL_UID (t)))
 	    {
-	      warning_at (OMP_CLAUSE_LOCATION (c), 0,
+	      warning_at (OMP_CLAUSE_LOCATION (c), OPT_Wopenmp,
 			  "%qE appears more than once in %<allocate%> clauses",
 			  t);
 	      remove = true;
@@ -16996,6 +16996,19 @@  c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	    }
 	  break;
 
+	case OMP_CLAUSE_INIT:
+	case OMP_CLAUSE_DESTROY:
+	case OMP_CLAUSE_USE:
+	  t = OMP_CLAUSE_DECL (c);
+	  if (bitmap_bit_p (&generic_head, DECL_UID (t)))
+	    {
+	      error_at (OMP_CLAUSE_LOCATION (c),
+			"%qD appears more than once in action clauses", t);
+	      remove = true;
+	    }
+	  bitmap_set_bit (&generic_head, DECL_UID (t));
+	  pc = &OMP_CLAUSE_CHAIN (c);
+	  break;
 	default:
 	  gcc_unreachable ();
 	}
diff --git a/gcc/fortran/dump-parse-tree.cc b/gcc/fortran/dump-parse-tree.cc
index ea5d2ab66a6..f749104dc5d 100644
--- a/gcc/fortran/dump-parse-tree.cc
+++ b/gcc/fortran/dump-parse-tree.cc
@@ -1544,63 +1544,42 @@  show_omp_namelist (int list_type, gfc_omp_namelist *n)
 	    fputs ("target,", dumpfile);
 	  if (n->u.init.targetsync)
 	    fputs ("targetsync,", dumpfile);
-	  if (n->u2.init_interop_fr)
+	  if (n->u2.init_interop)
 	    {
-	      char *attr_str = n->u.init.attr;
-	      int idx = 0;
-	      int fr_id;
+	      char *str = n->u2.init_interop;
 	      fputs ("prefer_type(", dumpfile);
-	      do
+	      while (str[0] == (char) GOMP_INTEROP_IFR_SEPARATOR)
 		{
-		  fr_id = n->u2.init_interop_fr[idx];
+		  bool has_fr = false;
 		  fputc ('{', dumpfile);
-		  if (fr_id != GOMP_INTEROP_IFR_NONE)
+		  str++;
+		  while (str[0] != (char) GOMP_INTEROP_IFR_SEPARATOR)
 		    {
-		      fputs ("fr(", dumpfile);
-		      do
-			{
-			  const char *fr_str = omp_get_name_from_fr_id (fr_id);
-			  if (fr_str)
-			    fprintf (dumpfile, "\"%s\"", fr_str);
-			  else
-			    fprintf (dumpfile, "%d", fr_id);
-			  fr_id = n->u2.init_interop_fr[++idx];
-			  if (fr_id != GOMP_INTEROP_IFR_SEPARATOR)
-			    fputc (',', dumpfile);
-			}
-		      while (fr_id != GOMP_INTEROP_IFR_SEPARATOR);
-		      fputc (')', dumpfile);
-		      if (attr_str && (attr_str[0] != ' ' || attr_str[1] != '\0'))
+		      if (has_fr)
 			fputc (',', dumpfile);
+		      has_fr = true;
+		      fputs ("fr(\"", dumpfile);
+		      fputs (omp_get_name_from_fr_id (str[0]), dumpfile);
+		      fputs ("\")", dumpfile);
+		      str++;
 		    }
-		  else
-		    fr_id = n->u2.init_interop_fr[++idx];
-		  if (attr_str && attr_str[0] == ' ' && attr_str[1] == '\0')
-		    attr_str += 2;
-		  else if (attr_str)
+		  str++;
+		  if (has_fr && str[0] != '\0')
+		    fputc (',', dumpfile);
+		  while (str[0] != '\0')
 		    {
 		      fputs ("attr(\"", dumpfile);
-		      do
-			{
-			  fputs ((char *) attr_str, dumpfile);
-			  fputc ('"', dumpfile);
-			  attr_str += strlen (attr_str) + 1;
-			  if (attr_str[0] == '\0')
-			    break;
-			  fputs (",\"", dumpfile);
-			}
-		      while (true);
-		      fputc (')', dumpfile);
+		      fputs (str, dumpfile);
+		      fputs ("\")", dumpfile);
+		      str += strlen (str) + 1;
+		      if (str[0] != '\0')
+			fputc (',', dumpfile);
 		    }
+		  str++;
 		  fputc ('}', dumpfile);
-		  fr_id = n->u2.init_interop_fr[++idx];
-		  if (fr_id == GOMP_INTEROP_IFR_SEPARATOR)
-		    break;
-		  fputc (',', dumpfile);
-		  if (attr_str)
-		    ++attr_str;
+		  if (str[0] != '\0')
+		    fputs (", ", dumpfile);
 		}
-	      while (true);
 	      fputc (')', dumpfile);
 	    }
 	  fputc (':', dumpfile);
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index 309095d74d5..d08439019a3 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -1403,7 +1403,6 @@  typedef struct gfc_omp_namelist
       bool present_modifier;
       struct
 	{
-	  char *attr;
 	  int len;
 	  bool target;
 	  bool targetsync;
@@ -1416,7 +1415,7 @@  typedef struct gfc_omp_namelist
       gfc_expr *allocator;
       struct gfc_symbol *traits_sym;
       struct gfc_omp_namelist *duplicate_of;
-      char *init_interop_fr;
+      char *init_interop;
     } u2;
   struct gfc_omp_namelist *next;
   locus where;
diff --git a/gcc/fortran/match.cc b/gcc/fortran/match.cc
index 6e9da8c3e68..75dd2693fe6 100644
--- a/gcc/fortran/match.cc
+++ b/gcc/fortran/match.cc
@@ -5608,7 +5608,7 @@  gfc_free_omp_namelist (gfc_omp_namelist *name, bool free_ns,
 {
   gfc_omp_namelist *n;
   gfc_expr *last_allocator = NULL;
-  char *last_init_attr = NULL;
+  char *last_init_interop = NULL;
 
   for (; name; name = n)
     {
@@ -5632,11 +5632,10 @@  gfc_free_omp_namelist (gfc_omp_namelist *name, bool free_ns,
 	{ }  /* name->u2.traits_sym: shall not call gfc_free_symbol here. */
       else if (free_init)
 	{
-	  if (name->u.init.attr != last_init_attr)
+	  if (name->u2.init_interop != last_init_interop)
 	    {
-	      last_init_attr = name->u.init.attr;
-	      free (name->u.init.attr);
-	      free (name->u2.init_interop_fr);
+	      last_init_interop = name->u2.init_interop;
+	      free (name->u2.init_interop);
 	    }
 	}
       else if (name->u2.udr)
diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index 9fb3836da31..8dc486a56a2 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -1835,45 +1835,42 @@  error:
      prefer_type ( <const-int-expr|string literal> [, ...]
    or
      prefer_type ( '{' <fr(...) | attr (...)>, ...] '}' [, '{' ... '}' ] )
-   where 'fr' takes an integer named constant or a string literal
-   and 'attr takes a string literal, starting with 'ompx_')
+   where 'fr' takes a constant expression or a string literal
+   and 'attr takes a list of string literals, starting with 'ompx_')
 
    For the foreign runtime identifiers, string values are converted to
-   their integer value; unknown string or integer values are set to 0.
-
-   For the simple syntax, pref_int_array contains alternatingly the
-   fr_id integer value and GOMP_INTEROP_IFR_SEPARATOR followed by a
-   GOMP_INTEROP_IFR_SEPARATOR as last item.
-   For the complex syntax, it contains the values associated with a
-   'fr(...)' followed by GOMP_INTEROP_IFR_SEPARATOR.  If there is no
-   'fr' in a curly-brace block, it is GOMP_INTEROP_IFR_NONE followed
-   by GOMP_INTEROP_IFR_SEPARATOR.  An additional GOMP_INTEROP_IFR_SEPARATOR
-   at the end terminates the array.
-
-   For attributes, if the simply syntax is used, it is NULL - likewise if no
-   'attr' appears.  For the complex syntax it is: For reach curly-brace block,
-   it is \0\0 is no attr appears and otherwise a concatenation (including
-   the \0) of all 'attr' strings followed by a tailing '\0'. At the end,
-   another '\0' follows.  */
+   their integer value; unknown string or integer values are set to
+   GOMP_INTEROP_IFR_KNOWN.
+
+   Data format:
+    For the foreign runtime identifiers, string values are converted to
+    their integer value; unknown string or integer values are set to 0.
+
+    Each item (a) GOMP_INTEROP_IFR_SEPARATOR
+	      (b) for any 'fr', its integer value.
+		  Note: Spec only permits 1 'fr' entry (6.0; changed after TR13)
+	      (c) GOMP_INTEROP_IFR_SEPARATOR
+	      (d) list of \0-terminated non-empty strings for 'attr'
+	      (e) '\0'
+    Tailing '\0'.  */
 
 static match
-gfc_match_omp_prefer_type (char **fr_int_array, char **attr_str, int *attr_str_len)
+gfc_match_omp_prefer_type (char **type_str, int *type_str_len)
 {
   gfc_expr *e;
-  int cnt_brace_grp = 0;
-  std::vector<char> int_list;
-  std::string attr_string;
+  std::string type_string, attr_string;
   /* New syntax.  */
   if (gfc_peek_ascii_char () == '{')
     do
       {
+	attr_string.clear ();
+	type_string += (char) GOMP_INTEROP_IFR_SEPARATOR;
 	if (gfc_match ("{ ") != MATCH_YES)
 	  {
 	    gfc_error ("Expected %<{%> at %C");
 	    return MATCH_ERROR;
 	  }
 	bool fr_found = false;
-	bool attr_found = false;
 	do
 	  {
 	    if (gfc_match ("fr ( ") == MATCH_YES)
@@ -1887,24 +1884,27 @@  gfc_match_omp_prefer_type (char **fr_int_array, char **attr_str, int *attr_str_l
 		fr_found = true;
 		do
 		  {
-		    if (gfc_match_expr (&e) != MATCH_YES)
-		      return MATCH_ERROR;
-		    if (e->expr_type != EXPR_CONSTANT
-			|| e->ref != NULL
+		    bool found_literal = false;
+		    match m = MATCH_YES;
+		    if (gfc_match_literal_constant (&e, false) == MATCH_YES)
+		      found_literal = true;
+		    else
+		      m = gfc_match_expr (&e);
+		    if (m != MATCH_YES
 			|| !gfc_resolve_expr (e)
+			|| e->rank != 0
+			|| e->expr_type != EXPR_CONSTANT
 			|| (e->ts.type != BT_INTEGER
-			    && e->ts.type != BT_CHARACTER)
+			    && (!found_literal || e->ts.type != BT_CHARACTER))
 			|| (e->ts.type == BT_INTEGER
-			    && (!e->symtree
-				|| e->symtree->n.sym->attr.flavor != FL_PARAMETER
-				|| !mpz_fits_sint_p (e->value.integer)))
+			    && !mpz_fits_sint_p (e->value.integer))
 			|| (e->ts.type == BT_CHARACTER
 			    && (e->ts.kind != gfc_default_character_kind
-				|| e->value.character.length == 0)))
+			|| e->value.character.length == 0)))
 		      {
-			gfc_error ("Expected scalar integer parameter or "
-				   "non-empty default-kind character literal "
-				   "at %L", &e->where);
+			gfc_error ("Expected constant scalar integer expression"
+				   " or non-empty default-kind character "
+				   "literal at %L", &e->where);
 			gfc_free_expr (e);
 			return MATCH_ERROR;
 		      }
@@ -1915,10 +1915,11 @@  gfc_match_omp_prefer_type (char **fr_int_array, char **attr_str, int *attr_str_l
 			val = mpz_get_si (e->value.integer);
 			if (val < 1 || val > GOMP_INTEROP_IFR_LAST)
 			  {
-			    gfc_warning (OPT_Wopenmp,
-					 "Unknown foreign runtime identifier "
-					 "%qd at %L", val, &e->where);
-			    val = 0;
+			    gfc_warning_now (OPT_Wopenmp,
+					     "Unknown foreign runtime "
+					     "identifier %qd at %L",
+					     val, &e->where);
+			    val = GOMP_INTEROP_IFR_UNKNOWN;
 			  }
 		      }
 		    else
@@ -1934,40 +1935,30 @@  gfc_match_omp_prefer_type (char **fr_int_array, char **attr_str, int *attr_str_l
 			    return MATCH_ERROR;
 			  }
 			val = omp_get_fr_id_from_name (str);
-			if (val == 0)
-			  gfc_warning (OPT_Wopenmp,
-				       "Unknown foreign runtime identifier %qs "
-				       "at %L", str, &e->where);
+			if (val == GOMP_INTEROP_IFR_UNKNOWN)
+			  gfc_warning_now (OPT_Wopenmp,
+					   "Unknown foreign runtime identifier "
+					   "%qs at %L", str, &e->where);
 		      }
-		    int_list.push_back (val);
-		    if (gfc_match (", ") == MATCH_YES)
-		      continue;
+
+		    type_string += (char) val;
 		    if (gfc_match (") ") == MATCH_YES)
 		      break;
-		    gfc_error ("Expected %<,%> or %<)%> at %C");
+		    gfc_error ("Expected %<)%> at %C");
 		    return MATCH_ERROR;
 		  }
 		while (true);
 	      }
 	    else if (gfc_match ("attr ( ") == MATCH_YES)
 	      {
-		attr_found = true;
-		if (attr_string.empty ())
-		  for (int i = 0; i < cnt_brace_grp; ++i)
-		    {
-		      /* Add dummy elements for previous curly-brace blocks.  */
-		      attr_string += ' ';
-		      attr_string += '\0';
-		      attr_string += '\0';
-		    }
 		do
 		  {
-		    if (gfc_match_expr (&e) != MATCH_YES)
-		      return MATCH_ERROR;
-		    if (e->expr_type != EXPR_CONSTANT
+		    if (gfc_match_literal_constant (&e, false) != MATCH_YES
+			|| !gfc_resolve_expr (e)
+			|| e->expr_type != EXPR_CONSTANT
 			|| e->rank != 0
 			|| e->ts.type != BT_CHARACTER
-		    || e->ts.kind != gfc_default_character_kind)
+			|| e->ts.kind != gfc_default_character_kind)
 		      {
 			gfc_error ("Expected default-kind character literal "
 				   "at %L", &e->where);
@@ -2016,21 +2007,9 @@  gfc_match_omp_prefer_type (char **fr_int_array, char **attr_str, int *attr_str_l
 	    return MATCH_ERROR;
 	  }
 	while (true);
-	++cnt_brace_grp;
-	if (!fr_found)
-	  int_list.push_back (GOMP_INTEROP_IFR_NONE);
-	int_list.push_back (GOMP_INTEROP_IFR_SEPARATOR);
-	if (!attr_string.empty ())
-	  {
-	    if (!attr_found)
-	      {
-		/* Dummy entry.  */
-		attr_string += ' ';
-		attr_string += '\0';
-	      }
-	    attr_string += '\0';
-	  }
-
+	type_string += (char) GOMP_INTEROP_IFR_SEPARATOR;
+	type_string += attr_string;
+	type_string += '\0';
 	if (gfc_match (", ") == MATCH_YES)
 	  continue;
 	if (gfc_match (") ") == MATCH_YES)
@@ -2042,12 +2021,19 @@  gfc_match_omp_prefer_type (char **fr_int_array, char **attr_str, int *attr_str_l
   else
     do
       {
-	if (gfc_match_expr (&e) != MATCH_YES)
-	  return MATCH_ERROR;
-	if (!gfc_resolve_expr (e)
+	type_string += (char) GOMP_INTEROP_IFR_SEPARATOR;
+	bool found_literal = false;
+	match m = MATCH_YES;
+	if (gfc_match_literal_constant (&e, false) == MATCH_YES)
+	  found_literal = true;
+	else
+	  m = gfc_match_expr (&e);
+	if (m != MATCH_YES
+	    || !gfc_resolve_expr (e)
 	    || e->rank != 0
 	    || e->expr_type != EXPR_CONSTANT
-	    || (e->ts.type != BT_INTEGER && e->ts.type != BT_CHARACTER)
+	    || (e->ts.type != BT_INTEGER
+		&& (!found_literal || e->ts.type != BT_CHARACTER))
 	    || (e->ts.type == BT_INTEGER
 		&& !mpz_fits_sint_p (e->value.integer))
 	    || (e->ts.type == BT_CHARACTER
@@ -2066,9 +2052,9 @@  gfc_match_omp_prefer_type (char **fr_int_array, char **attr_str, int *attr_str_l
 	    val = mpz_get_si (e->value.integer);
 	    if (val < 1 || val > GOMP_INTEROP_IFR_LAST)
 	      {
-		gfc_warning (OPT_Wopenmp,
-			     "Unknown foreign runtime identifier %qd at %L",
-			     val, &e->where);
+		gfc_warning_now (OPT_Wopenmp,
+				 "Unknown foreign runtime identifier %qd at %L",
+				 val, &e->where);
 		val = 0;
 	      }
 	  }
@@ -2084,13 +2070,14 @@  gfc_match_omp_prefer_type (char **fr_int_array, char **attr_str, int *attr_str_l
 		return MATCH_ERROR;
 	      }
 	    val = omp_get_fr_id_from_name (str);
-	    if (val == 0)
-	      gfc_warning (OPT_Wopenmp,
-			   "Unknown foreign runtime identifier %qs at %L",
-			   str, &e->where);
+	    if (val == GOMP_INTEROP_IFR_UNKNOWN)
+	      gfc_warning_now (OPT_Wopenmp,
+			       "Unknown foreign runtime identifier %qs at %L",
+			       str, &e->where);
 	  }
-	int_list.push_back (val);
-	int_list.push_back (GOMP_INTEROP_IFR_SEPARATOR);
+	type_string += (char) val;
+	type_string += (char) GOMP_INTEROP_IFR_SEPARATOR;
+	type_string += '\0';
 	gfc_free_expr (e);
 	if (gfc_match (", ") == MATCH_YES)
 	  continue;
@@ -2100,17 +2087,10 @@  gfc_match_omp_prefer_type (char **fr_int_array, char **attr_str, int *attr_str_l
 	return MATCH_ERROR;
       }
     while (true);
-  int_list.push_back (GOMP_INTEROP_IFR_SEPARATOR);
-  *fr_int_array = XNEWVEC (char, int_list.size ());
-  memcpy (*fr_int_array, int_list.data (), sizeof (char) * int_list.size ());
-
-  if (!attr_string.empty ())
-    {
-      attr_string += '\0';
-      *attr_str_len = attr_string.length();
-      *attr_str = XNEWVEC (char, attr_string.length ());
-      memcpy (*attr_str, attr_string.data (), attr_string.length ());
-    }
+  type_string += '\0';
+  *type_str_len = type_string.length();
+  *type_str = XNEWVEC (char, type_string.length ());
+  memcpy (*type_str, type_string.data (), type_string.length ());
   return MATCH_YES;
 }
 
@@ -2122,21 +2102,19 @@  static match
 gfc_match_omp_init (gfc_omp_namelist **list)
 {
   bool target = false, targetsync = false;
-  char *fr_int_array = NULL;
-  char *attr_str = NULL;
-  int attr_str_len = 0;
+  char *type_str = NULL;
+  int type_str_len = 0;
   match m;
   locus old_loc = gfc_current_locus;
   do {
        if (gfc_match ("prefer_type ( ") == MATCH_YES)
 	{
-	  if (fr_int_array)
+	  if (type_str)
 	    {
 	      gfc_error ("Duplicate %<prefer_type%> modifier at %C");
 	      return MATCH_ERROR;
 	    }
-	  m = gfc_match_omp_prefer_type (&fr_int_array, &attr_str,
-					 &attr_str_len);
+	  m = gfc_match_omp_prefer_type (&type_str, &type_str_len);
 	  if (m != MATCH_YES)
 	    return m;
 	  if (gfc_match (", ") == MATCH_YES)
@@ -2148,16 +2126,21 @@  gfc_match_omp_init (gfc_omp_namelist **list)
 	}
        if (gfc_match ("targetsync ") == MATCH_YES)
 	{
+	  if (targetsync)
+	    {
+	      /* Avoid the word 'modifier' as it could be also be no clauses and
+		 twice a variable named 'targetsync', which is also invalid.  */
+	      gfc_error ("Duplicate %<targetsync%> at %C");
+	      return MATCH_ERROR;
+	    }
 	  targetsync = true;
 	  if (gfc_match (", ") == MATCH_YES)
 	    continue;
 	  if (gfc_match (": ") == MATCH_YES)
 	    break;
 	  gfc_char_t c = gfc_peek_char ();
-	  if (!fr_int_array
-	      && (c == ')'
-		  || (gfc_current_form != FORM_FREE
-		      && (c == '_' || ISALPHA (c)))))
+	  if (!type_str && (c == ')' || (gfc_current_form != FORM_FREE
+					 && (c == '_' || ISALPHA (c)))))
 	    {
 	      gfc_current_locus = old_loc;
 	      break;
@@ -2167,16 +2150,19 @@  gfc_match_omp_init (gfc_omp_namelist **list)
 	}
       if (gfc_match ("target ") == MATCH_YES)
 	{
+	  if (target)
+	    {
+	      gfc_error ("Duplicate %<target%> at %C");
+	      return MATCH_ERROR;
+	    }
 	  target = true;
 	  if (gfc_match (", ") == MATCH_YES)
 	    continue;
 	  if (gfc_match (": ") == MATCH_YES)
 	    break;
 	  gfc_char_t c = gfc_peek_char ();
-	  if (!fr_int_array
-	      && (c == ')'
-		  || (gfc_current_form != FORM_FREE
-		      && (c == '_' || ISALPHA (c)))))
+	  if (!type_str && (c == ')' || (gfc_current_form != FORM_FREE
+					 && (c == '_' || ISALPHA (c)))))
 	    {
 	      gfc_current_locus = old_loc;
 	      break;
@@ -2184,7 +2170,7 @@  gfc_match_omp_init (gfc_omp_namelist **list)
 	  gfc_error ("Expected %<,%> or %<:%> at %C");
 	  return MATCH_ERROR;
 	}
-      if (fr_int_array)
+      if (type_str)
 	{
 	  gfc_error ("Expected %<target%> or %<targetsync%> at %C");
 	  return MATCH_ERROR;
@@ -2201,9 +2187,8 @@  gfc_match_omp_init (gfc_omp_namelist **list)
    {
      n->u.init.target = target;
      n->u.init.targetsync = targetsync;
-     n->u.init.attr = attr_str;
-     n->u.init.len = attr_str_len;
-     n->u2.init_interop_fr = fr_int_array;
+     n->u.init.len = type_str_len;
+     n->u2.init_interop = type_str;
    }
  return MATCH_YES;
 }
@@ -8459,6 +8444,21 @@  resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
 	    break;
 	  }
     }
+  if (code && code->op == EXEC_OMP_INTEROP)
+    for (list = OMP_LIST_INIT; list <= OMP_LIST_DESTROY; list++)
+      for (n = omp_clauses->lists[list]; n; n = n->next)
+	{
+	  if (n->sym->ts.type != BT_INTEGER
+	      || n->sym->ts.kind != gfc_index_integer_kind
+	      || n->sym->attr.dimension
+	      || n->sym->attr.flavor != FL_VARIABLE)
+	    gfc_error ("%qs at %L in %qs clause must be a scalar integer "
+		       "variable of %<omp_interop_kind%> kind", n->sym->name,
+		       &n->where, clause_names[list]);
+	  if (list != OMP_LIST_USE && n->sym->attr.intent == INTENT_IN)
+	    gfc_error ("%qs at %L in %qs clause must be definable",
+		       n->sym->name, &n->where, clause_names[list]);
+	}
 
   /* Detect specifically the case where we have "map(x) private(x)" and raise
      an error.  If we have "...simd" combined directives though, the "private"
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index 6c2c7482516..4f4b408d61c 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -2775,12 +2775,56 @@  gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 	case OMP_LIST_SCAN_EX:
 	  clause_code = OMP_CLAUSE_EXCLUSIVE;
 	  goto add_clause;
+	case OMP_LIST_USE:
+	  clause_code = OMP_CLAUSE_USE;
+	  goto add_clause;
+	case OMP_LIST_DESTROY:
+	  clause_code = OMP_CLAUSE_DESTROY;
+	  goto add_clause;
 
 	add_clause:
 	  omp_clauses
 	    = gfc_trans_omp_variable_list (clause_code, n, omp_clauses,
 					   declare_simd);
 	  break;
+
+	case OMP_LIST_INIT:
+	  {
+	    tree pref_type = NULL_TREE;
+	    const char *last = NULL;
+	    for (; n != NULL; n = n->next)
+	      if (n->sym->attr.referenced)
+		{
+		  tree t = gfc_trans_omp_variable (n->sym, false);
+		  if (t == error_mark_node)
+		    continue;
+		  tree node = build_omp_clause (input_location,
+						OMP_CLAUSE_INIT);
+		  OMP_CLAUSE_DECL (node) = t;
+		  if (n->u.init.target)
+		    OMP_CLAUSE_INIT_TARGET (node) = 1;
+		  if (n->u.init.targetsync)
+		    OMP_CLAUSE_INIT_TARGETSYNC (node) = 1;
+		  if (last != n->u2.init_interop)
+		    {
+		      last = n->u2.init_interop;
+		      if (n->u2.init_interop == NULL)
+			pref_type = NULL_TREE;
+		      else
+			{
+			  pref_type = build_string (n->u.init.len,
+						    n->u2.init_interop);
+			  TREE_TYPE (pref_type)
+			    = build_array_type_nelts (unsigned_char_type_node,
+						      n->u.init.len);
+			}
+		    }
+		  OMP_CLAUSE_INIT_PREFER_TYPE (node) = pref_type;
+		  omp_clauses = gfc_trans_add_clause (node, omp_clauses);
+		}
+	    break;
+	  }
+
 	case OMP_LIST_ALIGNED:
 	  for (; n != NULL; n = n->next)
 	    if (n->sym->attr.referenced || declare_simd)
@@ -8027,6 +8071,18 @@  gfc_trans_omp_target_update (gfc_code *code)
   return gfc_finish_block (&block);
 }
 
+static tree
+gfc_trans_openmp_interop (gfc_code *code, gfc_omp_clauses *clauses)
+{
+  stmtblock_t block;
+  gfc_start_block (&block);
+  tree omp_clauses = gfc_trans_omp_clauses (&block, clauses, code->loc);
+  tree stmt = build1_loc (input_location, OMP_INTEROP, void_type_node,
+			  omp_clauses);
+  gfc_add_expr_to_block (&block, stmt);
+  return gfc_finish_block (&block);
+}
+
 static tree
 gfc_trans_omp_workshare (gfc_code *code, gfc_omp_clauses *clauses)
 {
@@ -8365,8 +8421,7 @@  gfc_trans_omp_directive (gfc_code *code)
     case EXEC_OMP_WORKSHARE:
       return gfc_trans_omp_workshare (code, code->ext.omp_clauses);
     case EXEC_OMP_INTEROP:
-      sorry ("%<!$OMP INTEROP%>");
-      return build_empty_stmt (input_location);
+      return gfc_trans_openmp_interop (code, code->ext.omp_clauses);
     default:
       gcc_unreachable ();
     }
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 827941b24db..77259de5f12 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -18846,6 +18846,11 @@  gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
 	    break;
 	  }
 
+	case OMP_INTEROP:
+	  sorry_at (EXPR_LOCATION (*expr_p),
+		    "%<#pragma omp interop%> not yet supported");
+	  ret = GS_ERROR;
+	  break;
 	case OMP_ATOMIC:
 	case OMP_ATOMIC_READ:
 	case OMP_ATOMIC_CAPTURE_OLD:
diff --git a/gcc/omp-general.cc b/gcc/omp-general.cc
index 72fb7f92ff7..397db1b9a86 100644
--- a/gcc/omp-general.cc
+++ b/gcc/omp-general.cc
@@ -3500,7 +3500,7 @@  omp_get_fr_id_from_name (const char *str)
   for (unsigned i = 0; i < ARRAY_SIZE (omp_interop_fr_str); ++i)
     if (!strcmp (str, omp_interop_fr_str[i]))
       return i + 1;
-  return 0;
+  return GOMP_INTEROP_IFR_UNKNOWN;
 }
 
 /* Returns the string value to a foreign-runtime integer value or NULL if value
@@ -3510,7 +3510,7 @@  const char *
 omp_get_name_from_fr_id (int fr_id)
 {
   if (fr_id < 1 || fr_id > (int) ARRAY_SIZE (omp_interop_fr_str))
-    return NULL;
+    return "<unknown>";
   return omp_interop_fr_str[fr_id-1];
 }
 
diff --git a/gcc/testsuite/c-c++-common/gomp/interop-1.c b/gcc/testsuite/c-c++-common/gomp/interop-1.c
new file mode 100644
index 00000000000..71fdf484d0c
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/interop-1.c
@@ -0,0 +1,115 @@ 
+/* { dg-skip-if "C++ not yet implemented" { c++ } }  */
+
+/* Because of 'constexpr':  */
+/* { dg-additional-options "-std=c23"  { target c } } */
+
+/* { dg-prune-output "sorry, unimplemented: '#pragma omp interop' not yet supported" }  */
+
+/* The following definitions are in omp_lib, which cannot be included
+   in gcc/testsuite/  */
+
+#if __cplusplus >= 201103L
+# define __GOMP_UINTPTR_T_ENUM : omp_uintptr_t
+#else
+# define __GOMP_UINTPTR_T_ENUM
+#endif
+
+typedef enum omp_interop_t __GOMP_UINTPTR_T_ENUM
+{
+  omp_interop_none = 0,
+  __omp_interop_t_max__ = __UINTPTR_MAX__
+} omp_interop_t;
+
+typedef enum omp_interop_fr_t
+{
+  omp_ifr_cuda = 1,
+  omp_ifr_cuda_driver = 2,
+  omp_ifr_opencl = 3,
+  omp_ifr_sycl = 4,
+  omp_ifr_hip = 5,
+  omp_ifr_level_zero = 6,
+  omp_ifr_hsa = 7,
+  omp_ifr_last = omp_ifr_hsa
+} omp_interop_fr_t;
+
+
+// ---------------------------------
+
+void f()
+{
+constexpr omp_interop_fr_t ifr_scalar = omp_ifr_hsa;
+constexpr omp_interop_fr_t ifr_array[] = {omp_ifr_cuda, omp_ifr_hip};
+constexpr char my_string[] = "cuda";
+omp_interop_t obj1, obj2, obj3, obj4, obj5;
+int x;
+
+#pragma omp interop init(obj1) init(target,targetsync : obj2, obj3) nowait   // OK
+#pragma omp interop init(obj1) init (targetsync  : obj2, obj3) nowait   // OK
+#pragma omp interop init(obj1) init (targetsync , target : obj2, obj3) nowait   // OK
+
+#pragma omp interop init(obj1) init(target,targetsync,target: obj2, obj3) nowait   // { dg-error "duplicate 'target' modifier" }
+#pragma omp interop init(obj1) init(target,targetsync, targetsync : obj2, obj3) nowait   // { dg-error "duplicate 'targetsync' modifier" }
+
+#pragma omp interop init(prefer_type("cuda", omp_ifr_opencl, omp_ifr_level_zero, "hsa"), targetsync : obj1) \
+                    destroy(obj2, obj3) depend(inout: x) use(obj4, obj5) device(device_num: 0)
+
+#pragma omp interop init(prefer_type("cu" "da"), targetsync : obj1)   // OK
+
+#pragma omp assume contains(interop)
+  {
+  #pragma omp interop init(prefer_type("cuða") : obj3)  // { dg-warning "unknown foreign runtime identifier 'cu\[^'\]*a'" }
+  }
+
+#pragma omp interop init(prefer_type("cu\0da") : obj3) // { dg-error "string literal must not contain '\\\\0'" }
+
+#pragma omp interop depend(inout: x) , use(obj2), destroy(obj3) //  OK, use or destory might have 'targetsync'
+
+#pragma omp interop depend(inout: x) use(obj2), destroy(obj3) //  Likewise
+
+#pragma omp interop depend(inout: x) init(targetsync : obj5)  use(obj2), destroy(obj3) init(prefer_type("cuda"), targetsync : obj4) // OK
+
+#pragma omp interop init ( target , prefer_type( { fr("hsa") }, "hip") : obj1) // { dg-error "expected '\{' before string constant" }
+
+#pragma omp interop init ( target , prefer_type( { fr("hsa"), attr("ompx_nothing") , fr("hsa" ) }) :obj1) // { dg-error "duplicated 'fr' preference selector before '\\(' token" }
+
+#pragma omp interop init ( prefer_type( 4, omp_ifr_hip*4) : obj1)  // { dg-warning "unknown foreign runtime identifier '20'" }
+#pragma omp interop init ( prefer_type( __builtin_sin(3.3) : obj1)  // { dg-error "'prefer_type' undeclared \\(first use in this function\\)" }
+                         // { dg-error "expected '\\)' before '\\(' token" "" { target *-*-* } .-1 }
+#pragma omp interop init ( prefer_type( __builtin_sin(3.3) ) : obj1)  // { dg-error "expected string literal or constant integer expression before '\\)' token" }
+#pragma omp interop init ( prefer_type( {fr(4 ) }) : obj1) // OK
+#pragma omp interop init ( prefer_type( {fr("cu\0da" ) }) : obj1) // { dg-error "string literal must not contain '\\\\0'" }
+#pragma omp interop init ( prefer_type( {fr("cuda\0") }) : obj1) // { dg-error "string literal must not contain '\\\\0'" }
+#pragma omp interop init ( prefer_type( {fr("cuda" ) }) : obj1) // OK
+#pragma omp interop init ( prefer_type( {fr(omp_ifr_level_zero ) }, {fr(omp_ifr_hip)}) : obj1) // OK
+#pragma omp interop init ( prefer_type( {fr("cuda",  "cuda_driver") }) : obj1) // { dg-error "51: expected '\\)' before ',' token" }
+#pragma omp interop init ( prefer_type( {fr(my_string) }) : obj1) // { dg-error "54: expected string literal or constant integer expression before '\\)' token" }
+#pragma omp interop init ( prefer_type( {fr("hello" }) : obj1) // { dg-error "expected '\\)' before '\\(' token" }
+#pragma omp interop init ( prefer_type( {fr("hello") }) : obj1)
+/* { dg-warning "unknown foreign runtime identifier 'hello' \\\[-Wopenmp\\\]" "" { target *-*-* } .-1 }  */
+
+#pragma omp interop init ( prefer_type( {fr(x) }) : obj1) // { dg-error "expected string literal or constant integer expression before '\\)' token" }
+
+#pragma omp interop init ( prefer_type( {fr(ifr_scalar ) }) : obj1) // OK
+#pragma omp interop init ( prefer_type( {fr(ifr_array ) }) : obj1) // { dg-error "expected string literal or constant integer expression before '\\)' token" }
+// constexpr arrays are not part of C23; however, they are/were under consideration for C2y.
+#pragma omp interop init ( prefer_type( {fr(ifr_array[0] ) }) : obj1)  // { dg-error "expected string literal or constant integer expression before '\\)' token" }
+
+#pragma omp interop init ( prefer_type( omp_ifr_level_zero, omp_ifr_hip ) : obj1) // OK
+#pragma omp interop init ( prefer_type( omp_ifr_level_zero +1 ) : obj1) // OK
+#pragma omp interop init ( prefer_type( x ) : obj1) // { dg-error "expected string literal or constant integer expression before '\\)' token" }
+
+#pragma omp interop init ( prefer_type( ifr_scalar ) : obj1) // OK
+#pragma omp interop init ( prefer_type( ifr_array ) : obj1) // { dg-error "expected string literal or constant integer expression before '\\)' token" }
+// Unclear if okay:
+#pragma omp interop init ( prefer_type( ifr_array[1] ) : obj1) // { dg-error "expected string literal or constant integer expression before '\\)' token" }
+
+#pragma omp interop init ( prefer_type( 4, omp_ifr_hip*4) : obj1) // { dg-warning "unknown foreign runtime identifier '20'" }
+#pragma omp interop init ( prefer_type( 4, 1, 3) : obj1)
+
+#pragma omp interop init ( prefer_type( {fr("cuda") }, {fr(omp_ifr_hsa)} , {attr("ompx_a") } , {fr(omp_ifr_hip) }) : obj1)
+#pragma omp interop init ( prefer_type( {fr("cuda") }, {fr(omp_ifr_hsa,omp_ifr_level_zero)} , {attr("ompx_a") } , {fr(omp_ifr_hip) }) : obj1) // { dg-error "71: expected '\\)' before ',' token" }
+#pragma omp interop init ( prefer_type( {fr("cuda",5) }, {fr(omp_ifr_hsa,omp_ifr_level_zero)} , {attr("ompx_a") } , {fr(omp_ifr_hip) }) : obj1) // { dg-error "51: expected '\\)' before ',' token" }
+#pragma omp interop init ( prefer_type( {fr("sycl"), attr("ompx_1", "ompx_2"), attr("ompx_3") }, {attr("ompx_4", "ompx_5"),fr(omp_ifr_level_zero)} ) : obj1)
+#pragma omp interop init ( prefer_type( { fr(5), attr("ompx_1") }, {fr(omp_ifr_hsa)} , {attr("ompx_a") } ) : obj1)
+
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/interop-2.c b/gcc/testsuite/c-c++-common/gomp/interop-2.c
new file mode 100644
index 00000000000..96282a1875a
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/interop-2.c
@@ -0,0 +1,106 @@ 
+/* { dg-skip-if "C++ not yet implemented" { c++ } }  */
+
+/* Because of 'constexpr':  */
+/* { dg-additional-options "-std=c23"  { target c } } */
+
+/* { dg-prune-output "sorry, unimplemented: '#pragma omp interop' not yet supported" }  */
+
+/* The following definitions are in omp_lib, which cannot be included
+   in gcc/testsuite/  */
+
+#if __cplusplus >= 201103L
+# define __GOMP_UINTPTR_T_ENUM : omp_uintptr_t
+#else
+# define __GOMP_UINTPTR_T_ENUM
+#endif
+
+typedef enum omp_interop_t __GOMP_UINTPTR_T_ENUM
+{
+  omp_interop_none = 0,
+  __omp_interop_t_max__ = __UINTPTR_MAX__
+} omp_interop_t;
+
+typedef enum omp_interop_fr_t
+{
+  omp_ifr_cuda = 1,
+  omp_ifr_cuda_driver = 2,
+  omp_ifr_opencl = 3,
+  omp_ifr_sycl = 4,
+  omp_ifr_hip = 5,
+  omp_ifr_level_zero = 6,
+  omp_ifr_hsa = 7,
+  omp_ifr_last = omp_ifr_hsa
+} omp_interop_fr_t;
+
+
+// ---------------------------------
+
+void f(const omp_interop_t ocp)
+{
+  constexpr omp_interop_t oce = omp_interop_none;
+  const omp_interop_t occ = omp_interop_none;
+  omp_interop_t od[5];
+  omp_interop_t *op;
+  short o2;
+  float of;
+
+  #pragma omp interop init (ocp)  // { dg-error "'ocp' shall not be const" }
+  #pragma omp interop init (oce)  // { dg-error "'oce' shall not be const" }
+  #pragma omp interop init (occ)  // { dg-error "'occ' shall not be const" }
+  #pragma omp interop init (od)   // { dg-error "'od' must be of 'omp_interop_t'" }
+  #pragma omp interop init (od[1])// { dg-error "expected '\\)' before '\\\[' token" }
+                                  // { dg-error "'od' must be of 'omp_interop_t'" "" { target *-*-* } .-1 }
+  #pragma omp interop init (op)   // { dg-error "'op' must be of 'omp_interop_t'" }
+  #pragma omp interop init (*op)  // { dg-error "expected identifier before '\\*' token" }
+  #pragma omp interop init (o2)   // { dg-error "'o2' must be of 'omp_interop_t'" }
+  #pragma omp interop init (of)   // { dg-error "'of' must be of 'omp_interop_t'" }
+
+  #pragma omp interop use (ocp)  // OK
+  #pragma omp interop use (oce)  // odd but okay
+  #pragma omp interop use (occ)  // okayish
+  #pragma omp interop use (od)   // { dg-error "'od' must be of 'omp_interop_t'" }
+  #pragma omp interop use (od[1])// { dg-error "expected '\\)' before '\\\[' token" }
+                                 // { dg-error "'od' must be of 'omp_interop_t'" "" { target *-*-* } .-1 }
+  #pragma omp interop use (op)   // { dg-error "'op' must be of 'omp_interop_t'" }
+  #pragma omp interop use (*op)  // { dg-error "expected identifier before '\\*' token" }
+  #pragma omp interop use (o2)   // { dg-error "'o2' must be of 'omp_interop_t'" }
+  #pragma omp interop use (of)   // { dg-error "'of' must be of 'omp_interop_t'" }
+
+  #pragma omp interop destroy (ocp)  // { dg-error "'ocp' shall not be const" }
+  #pragma omp interop destroy (oce)  // { dg-error "'oce' shall not be const" }
+  #pragma omp interop destroy (occ)  // { dg-error "'occ' shall not be const" }
+  #pragma omp interop destroy (od)   // { dg-error "'od' must be of 'omp_interop_t'" }
+  #pragma omp interop destroy (od[1])// { dg-error "expected '\\)' before '\\\[' token" }
+                                     // { dg-error "'od' must be of 'omp_interop_t'" "" { target *-*-* } .-1 }
+  #pragma omp interop destroy (op)   // { dg-error "'op' must be of 'omp_interop_t'" }
+  #pragma omp interop destroy (*op)  // { dg-error "expected identifier before '\\*' token" }
+  #pragma omp interop destroy (o2)   // { dg-error "'o2' must be of 'omp_interop_t'" }
+  #pragma omp interop destroy (of)   // { dg-error "'of' must be of 'omp_interop_t'" }
+}
+
+void g()
+{
+  omp_interop_t obj1, obj2, obj3, obj4, obj5;
+  int x;
+
+  #pragma omp interop init ( prefer_type( {fr("") }) : obj1)  // { dg-error "non-empty string literal expected before '\\)' token" }
+  #pragma omp interop init ( prefer_type( {fr("hip") , attr(omp_ifr_cuda) }) : obj1) ! { dg-error "expected string literal before 'omp_ifr_cuda'" }
+
+  #pragma omp interop init ( prefer_type( {fr("hip") , attr("myooption") }) : obj1)  // { dg-error "'attr' string literal must start with 'ompx_'" }
+  #pragma omp interop init ( prefer_type( {fr("hip") , attr("ompx_option") , attr("ompx_") } ) : obj1)
+  #pragma omp interop init ( prefer_type( {fr("hip") , attr("ompx_option") }, { attr("ompx_") } ) : obj1)
+  #pragma omp interop init ( prefer_type( {fr("hip") , attr("ompx_option") }  { attr("ompx_") } ) : obj1)  // { dg-error "expected '\\)' or ',' before '\{' token" }
+  #pragma omp interop init ( prefer_type( {fr("hip") , attr("ompx_option")   ) : obj1)  // { dg-error "expected ',' or '\}' before '\\)' token" }
+
+  #pragma omp interop init ( prefer_type( {fr("hip") attr("ompx_option")   ) : obj1)  // { dg-error "expected ',' or '\}' before 'attr" }
+  #pragma omp interop init ( prefer_type( {fr("hip")}), prefer_type("cuda") : obj1)  // { dg-error "duplicate 'prefer_type' modifier" }
+
+  #pragma omp interop init ( prefer_type( {attr("ompx_option1,ompx_option2")   ) : obj1)  // { dg-error "'attr' string literal must not contain a comma" }
+
+  #pragma omp interop init ( targetsync other ) : obj1)   // { dg-error "'targetsync' undeclared \\(first use in this function\\)" }
+     // { dg-error "expected '\\)' before 'other'" "" { target *-*-*} .-1 }
+     // { dg-error "expected an OpenMP clause before ':' token" "" { target *-*-*} .-2 }
+  #pragma omp interop init ( prefer_type( {fr("cuda") } ), other : obj1)   // { dg-error "'init' clause with modifier other than 'prefer_type', 'target' or 'targetsync' before 'other'" }
+  #pragma omp interop init ( prefer_type( {fr("cuda") } ), obj1)   // { dg-error "'prefer_type' undeclared \\(first use in this function\\)" }
+                                                                   // { dg-error "expected '\\)' before '\\(' token" "" { target *-*-* } .-1 }
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/interop-3.c b/gcc/testsuite/c-c++-common/gomp/interop-3.c
new file mode 100644
index 00000000000..09866ff7b29
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/interop-3.c
@@ -0,0 +1,81 @@ 
+/* { dg-skip-if "C++ not yet implemented" { c++ } }  */
+
+/* { dg-prune-output "sorry, unimplemented: '#pragma omp interop' not yet supported" }  */
+
+/* The following definitions are in omp_lib, which cannot be included
+   in gcc/testsuite/  */
+
+#if __cplusplus >= 201103L
+# define __GOMP_UINTPTR_T_ENUM : omp_uintptr_t
+#else
+# define __GOMP_UINTPTR_T_ENUM
+#endif
+
+typedef enum omp_interop_t __GOMP_UINTPTR_T_ENUM
+{
+  omp_interop_none = 0,
+  __omp_interop_t_max__ = __UINTPTR_MAX__
+} omp_interop_t;
+
+typedef enum omp_interop_fr_t
+{
+  omp_ifr_cuda = 1,
+  omp_ifr_cuda_driver = 2,
+  omp_ifr_opencl = 3,
+  omp_ifr_sycl = 4,
+  omp_ifr_hip = 5,
+  omp_ifr_level_zero = 6,
+  omp_ifr_hsa = 7,
+  omp_ifr_last = omp_ifr_hsa
+} omp_interop_fr_t;
+
+
+// ---------------------------------
+
+void f()
+{
+  omp_interop_t obj1, obj2, obj3, obj4, obj5;
+  omp_interop_t target, targetsync, prefer_type;
+  int x;
+
+  #pragma omp interop init(obj1) init(target,targetsync : obj2, obj3) nowait
+
+  #pragma omp interop init(prefer_type("cuda", omp_ifr_opencl, omp_ifr_level_zero, "hsa"), targetsync : obj1) \
+                      destroy(obj2, obj3) depend(inout: x) use(obj4, obj5) device(device_num: 0)
+
+  #pragma omp assume contains(interop)
+    {
+    #pragma omp interop init(prefer_type("cu da") : obj3)  // { dg-warning "unknown foreign runtime identifier 'cu da'" }
+    }
+
+  #pragma omp interop init(obj1, obj2, obj1), use(obj4) destroy(obj4)
+  // { dg-error "'obj4' appears more than once in action clauses" "" { target *-*-* } .-1 }
+  // { dg-error "'obj1' appears more than once in action clauses" "" { target *-*-* } .-2 }
+
+  #pragma omp interop depend(inout: x)  // { dg-error "'depend' clause requires action clauses with 'targetsync' interop-type" }
+
+  #pragma omp interop depend(inout: x) , use(obj2), destroy(obj3) //  OK, use or destory might have 'targetsync'
+
+  #pragma omp interop depend(inout: x) use(obj2), destroy(obj3) //  Likewise
+
+  #pragma omp interop depend(inout: x) use(obj2), destroy(obj3) init(obj4) // { dg-error "'depend' clause requires action clauses with 'targetsync' interop-type" }
+  // { dg-note "69: 'init' clause lacks the 'targetsync' modifier" "" { target *-*-* } .-1 }
+
+  #pragma omp interop depend(inout: x) init(targetsync : obj5)  use(obj2), destroy(obj3) init(obj4) // { dg-error "'depend' clause requires action clauses with 'targetsync' interop-type" }
+  // { dg-note "'init' clause lacks the 'targetsync' modifier" "" { target *-*-* } .-1 }
+  #pragma omp interop depend(inout: x) init(targetsync : obj5)  use(obj2), destroy(obj3) init(prefer_type("cuda"), targetsync : obj4) // OK
+
+  #pragma omp interop init(target, targetsync, prefer_type, obj1)
+  #pragma omp interop init(prefer_type, obj1, target, targetsync)
+
+// Duplicated variable name or duplicated modifier:
+  #pragma omp interop init(target, targetsync,target : obj1)  // { dg-error "duplicate 'target' modifier" }
+  #pragma omp interop init(target, targetsync,target)         // { dg-error "'target' appears more than once in action clauses" }
+  #pragma omp interop init(target : target, targetsync,target)  // { dg-error "'target' appears more than once in action clauses" }
+
+  #pragma omp interop init(target, targetsync,targetsync : obj1)  // { dg-error "duplicate 'targetsync' modifier" }
+  #pragma omp interop init(target, targetsync,targetsync)         // { dg-error "targetsync' appears more than once in action clause" }
+  #pragma omp interop init(target : target, targetsync,targetsync)  // { dg-error "targetsync' appears more than once in action clause" }
+
+  #pragma omp interop init(, targetsync, prefer_type, obj1, target)  // { dg-error "expected identifier before ',' token" }
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/interop-4.c b/gcc/testsuite/c-c++-common/gomp/interop-4.c
new file mode 100644
index 00000000000..3b4939b3a43
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/interop-4.c
@@ -0,0 +1,76 @@ 
+/* { dg-skip-if "C++ not yet implemented" { c++ } }  */
+/* { dg-additional-options "-fdump-tree-original" }  */
+
+/* The following definitions are in omp_lib, which cannot be included
+   in gcc/testsuite/  */
+
+#if __cplusplus >= 201103L
+# define __GOMP_UINTPTR_T_ENUM : omp_uintptr_t
+#else
+# define __GOMP_UINTPTR_T_ENUM
+#endif
+
+typedef enum omp_interop_t __GOMP_UINTPTR_T_ENUM
+{ 
+  omp_interop_none = 0,
+  __omp_interop_t_max__ = __UINTPTR_MAX__
+} omp_interop_t;
+  
+typedef enum omp_interop_fr_t
+{ 
+  omp_ifr_cuda = 1,
+  omp_ifr_cuda_driver = 2,
+  omp_ifr_opencl = 3,
+  omp_ifr_sycl = 4,
+  omp_ifr_hip = 5,
+  omp_ifr_level_zero = 6,
+  omp_ifr_hsa = 7,
+  omp_ifr_last = omp_ifr_hsa
+} omp_interop_fr_t;
+
+void
+f()
+{
+  omp_interop_t obj1, obj2, obj3, obj4, obj5, obj6, obj7;
+  int x[6];
+
+  #pragma omp interop init ( obj1, obj2) use (obj3) destroy(obj4) init(obj5) destroy(obj6) use(obj7)   /* { dg-message "'#pragma omp interop' not yet supported" }  */
+  /* { dg-final { scan-tree-dump-times "#pragma omp interop use\\(obj7\\) destroy\\(obj6\\) init\\(obj5\\) destroy\\(obj4\\) use\\(obj3\\) init\\(obj2\\) init\\(obj1\\)\[\r\n\]" 1 "original" } }  */
+
+  #pragma omp interop nowait init (targetsync : obj1, obj2) use (obj3) destroy(obj4) init(target, targetsync : obj5) destroy(obj6) use(obj7) depend(inout: x)  /* { dg-message "'#pragma omp interop' not yet supported" }  */
+  /* { dg-final { scan-tree-dump-times "#pragma omp interop depend\\(inout:x\\) use\\(obj7\\) destroy\\(obj6\\) init\\(target, targetsync: obj5\\) destroy\\(obj4\\) use\\(obj3\\) init\\(targetsync: obj2\\) init\\(targetsync: obj1\\) nowait\[\r\n\]" 1 "original" } }  */
+
+  #pragma omp interop init ( obj1, obj2) init (target: obj3) init(targetsync : obj4) init(target,targetsync: obj5)  /* { dg-message "'#pragma omp interop' not yet supported" }  */
+  /* { dg-final { scan-tree-dump-times "#pragma omp interop init\\(target, targetsync: obj5\\) init\\(targetsync: obj4\\) init\\(target: obj3\\) init\\(obj2\\) init\\(obj1\\)\[\r\n\]" 1 "original" } }  */
+
+  /* --------------------------------------------  */
+
+  #pragma omp interop init (target, prefer_type(omp_ifr_cuda, omp_ifr_cuda+1, "hsa", "myPrivateInterop", omp_ifr_cuda-2) : obj1, obj2) init (target: obj3) init(prefer_type(omp_ifr_hip, "sycl", omp_ifr_opencl), targetsync : obj4, obj7) init(target,prefer_type("level_zero", omp_ifr_level_zero+0),targetsync: obj5)  /* { dg-message "'#pragma omp interop' not yet supported" }  */
+  /*
+     { dg-warning "unknown foreign runtime identifier 'myPrivateInterop' \\\[-Wopenmp\\\]" "" { target *-*-* } .-2 }
+     { dg-warning "unknown foreign runtime identifier '-1' \\\[-Wopenmp\\\]" "" { target *-*-* } .-3 }
+
+     { dg!final { scan-tree-dump-times "#pragma omp interop init\\(prefer_type\\({fr\\(\"level_zero\"\\)}, {fr\\(\"level_zero\"\\)}\\), target, targetsync: obj5\\) init\\(prefer_type\\({fr\\(\"hip\"\\)}, {fr\\(\"sycl\"\\)}, {fr\\(\"opencl\"\\)}\\), targetsync: obj7\\) init\\(prefer_type\\({fr\\(\"hip\"\\)}, {fr\\(\"sycl\"\\)}, {fr\\(\"opencl\"\\)}\\), targetsync: obj4\\) init\\(target: obj3\\) init\\(prefer_type\\({fr\\(\"cuda\"\\)}, {fr\\(\"cuda_driver\"\\)}, {fr\\(\"hsa\"\\)}, {fr\\(\"<unknown>\"\\)}, {fr\\(\"<unknown>\"\\)}\\), target: obj2\\) init\\(prefer_type\\({fr\\(\"cuda\"\\)}, {fr\\(\"cuda_driver\"\\)}, {fr\\(\"hsa\"\\)}, {fr\\(\"<unknown>\"\\)}, {fr\\(\"<unknown>\"\\)}\\), target: obj1\\)\[\r\n\]" 1 "original" } }
+  */
+
+
+/* -------------------------------------------- */
+
+  #pragma omp interop init ( target, prefer_type( {fr("hip"), attr("ompx_gnu_prio:1", "ompx_gnu_debug")}, {attr("ompx_gnu_nicest"), attr("ompx_something")}) : obj1, obj2) init ( prefer_type( {fr("cuda")}, {fr(omp_ifr_cuda_driver), attr("ompx_nix")}, {fr("best")}), targetsync : obj3, obj4) nowait use(obj5)  /* { dg-message "'#pragma omp interop' not yet supported" }  */
+  /*
+     { dg-warning "unknown foreign runtime identifier 'best' \\\[-Wopenmp\\\]" "" { target *-*-* } .-2 }
+
+     { dg-final { scan-tree-dump-times "#pragma omp interop use\\(obj5\\) nowait init\\(prefer_type\\({fr\\(\"cuda\"\\)}, {fr\\(\"cuda_driver\"\\),attr\\(\"ompx_nix\"\\)}, {fr\\(\"<unknown>\"\\)}\\), targetsync: obj4\\) init\\(prefer_type\\({fr\\(\"cuda\"\\)}, {fr\\(\"cuda_driver\"\\),attr\\(\"ompx_nix\"\\)}, {fr\\(\"<unknown>\"\\)}\\), targetsync: obj3\\) init\\(prefer_type\\({fr\\(\"hip\"\\),attr\\(\"ompx_gnu_prio:1\"\\),attr\\(\"ompx_gnu_debug\"\\)}, {attr\\(\"ompx_gnu_nicest\"\\),attr\\(\"ompx_something\"\\)}\\), target: obj2\\) init\\(prefer_type\\({fr\\(\"hip\"\\),attr\\(\"ompx_gnu_prio:1\"\\),attr\\(\"ompx_gnu_debug\"\\)}, {attr\\(\"ompx_gnu_nicest\"\\),attr\\(\"ompx_something\"\\)}\\), target: obj1\\)\[\r\n\]" 1 "original" } }
+  */
+
+}
+
+void
+g (int *y)
+{
+  omp_interop_t io1, io2, io3, io4, io5;
+
+  [[omp::directive (interop,init(prefer_type({fr("level_zero")}, {fr(omp_ifr_sycl),attr("ompx_in_order"),attr("ompx_queue:in_order")}), targetsync : io1, io2),use(io3),destroy(io4,io5),depend(inout:y),nowait)]];  /* { dg-message "'#pragma omp interop' not yet supported" }  */
+
+  /* { dg-final { scan-tree-dump-times "#pragma omp interop nowait depend\\(inout:y\\) destroy\\(io5\\) destroy\\(io4\\) use\\(io3\\) init\\(prefer_type\\(\{fr\\(\"level_zero\"\\)\}, \{fr\\(\"sycl\"\\),attr\\(\"ompx_in_order\"\\),attr\\(\"ompx_queue:in_order\"\\)\}\\), targetsync: io2\\) init\\(prefer_type\\(\{fr\\(\"level_zero\"\\)\}, \{fr\\(\"sycl\"\\),attr\\(\"ompx_in_order\"\\),attr\\(\"ompx_queue:in_order\"\\)\}\\), targetsync: io1\\)\[\r\n\]" 1 "original" } }  */
+}
diff --git a/gcc/testsuite/gfortran.dg/gomp/interop-1.f90 b/gcc/testsuite/gfortran.dg/gomp/interop-1.f90
index b7d2164812c..3d0cc9dd657 100644
--- a/gcc/testsuite/gfortran.dg/gomp/interop-1.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/interop-1.f90
@@ -33,16 +33,25 @@  integer(omp_interop_fr_kind), parameter :: ifr_array(2) = [omp_ifr_cuda, omp_ifr
 integer(omp_interop_kind) :: obj1, obj2, obj3, obj4, obj5
 integer :: x
 
-!$omp interop init(obj1) init(target,targetsync,target,targetsync : obj2, obj3) nowait
+!$omp interop init(obj1) init(target,targetsync : obj2, obj3) nowait ! OK
+!$omp interop init(obj1) init (targetsync  : obj2, obj3) nowait ! OK
+!$omp interop init(obj1) init (targetsync , target : obj2, obj3) nowait ! OK
 
-!$omp interop init(prefer_type("cu"//"da", omp_ifr_opencl, omp_ifr_level_zero, "hsa"), targetsync : obj1) &
+!$omp interop init(obj1) init(target,targetsync,target: obj2, obj3) nowait ! { dg-error "Duplicate 'target'" }
+!$omp interop init(obj1) init(target,targetsync, targetsync : obj2, obj3) nowait ! { dg-error "Duplicate 'targetsync'" }
+
+!$omp interop init(prefer_type("cuda", omp_ifr_opencl, omp_ifr_level_zero, "hsa"), targetsync : obj1) &
 !$omp&        destroy(obj2, obj3) depend(inout: x) use(obj4, obj5) device(device_num: 0)
 
+!$omp interop init(prefer_type("cu" // "da"), targetsync : obj1) ! { dg-error "37: Expected ',' or '\\)'" }
+! { dg-warning "Unknown foreign runtime identifier 'cu' at \\(1\\) \\\[-Wopenmp\\\]" "" { target *-*-* } .-1 }
+
 !$omp assume contains(interop)
-  !$omp interop init(prefer_type("cu"//char(1)//"da") : obj3)  ! { dg-warning "Unknown foreign runtime identifier 'cu\\\\x01da'" }
+  !$omp interop init(prefer_type("cuða") : obj3)  ! { dg-warning "Unknown foreign runtime identifier 'cu\[^'\]*a'" }
 !$omp end assume
 
-!$omp interop init(prefer_type("cu"//char(0)//"da") : obj3) ! { dg-error "Unexpected null character in character literal" }
+!$omp interop init(prefer_type("cu"//char(0)//"da") : obj3) ! { dg-error "36: Expected ',' or '\\)'" }
+! { dg-warning "Unknown foreign runtime identifier 'cu' at \\(1\\) \\\[-Wopenmp\\\]" "" { target *-*-* } .-1 }
 
 !$omp interop depend(inout: x) , use(obj2), destroy(obj3) !  OK, use or destory might have 'targetsync'
 
@@ -56,15 +65,19 @@  integer :: x
 
 !$omp interop init ( prefer_type( 4, omp_ifr_hip*4) : obj1)  ! { dg-warning "Unknown foreign runtime identifier '20'" }
 !$omp interop init ( prefer_type( sin(3.3) : obj1)  ! { dg-error "Expected constant scalar integer expression or non-empty default-kind character literal" }
-!$omp interop init ( prefer_type( {fr(4 ) }) : obj1) ! { dg-error "Expected scalar integer parameter or non-empty default-kind character literal" }
-!$omp interop init ( prefer_type( {fr(4_"cuda" ) }) : obj1) ! { dg-error "Expected scalar integer parameter or non-empty default-kind character literal" }
+!$omp interop init ( prefer_type( {fr(4 ) }) : obj1) ! OK
+!$omp interop init ( prefer_type( {fr(4_"cuda" ) }) : obj1) ! { dg-error "Expected constant scalar integer expression or non-empty default-kind character literal" }
 !$omp interop init ( prefer_type( {fr(c_char_"cuda") }) : obj1) ! OK
 !$omp interop init ( prefer_type( {fr(1_"cuda" ) }) : obj1) ! OK
 !$omp interop init ( prefer_type( {fr(omp_ifr_level_zero ) }, {fr(omp_ifr_hip)}) : obj1) ! OK
-!$omp interop init ( prefer_type( {fr(omp_ifr_level_zero + 1) }) : obj1) ! { dg-error "Expected scalar integer parameter or non-empty default-kind character literal" }
-!$omp interop init ( prefer_type( {fr(x) }) : obj1) ! { dg-error "Expected scalar integer parameter or non-empty default-kind character literal" }
-!$omp interop init ( prefer_type( {fr(ifr_array ) }) : obj1) ! { dg-error "Expected scalar integer parameter or non-empty default-kind character literal" }
-!$omp interop init ( prefer_type( {fr(ifr_array(1) ) }) : obj1) ! { dg-error "Expected scalar integer parameter or non-empty default-kind character literal" }
+!$omp interop init ( prefer_type( {fr("cuda" // "_driver") }) : obj1) ! { dg-error "46: Expected '\\)'" }
+!$omp interop init ( prefer_type( {fr(trim("cuda" // "_driver")) }) : obj1) ! { dg-error "38: Expected constant scalar integer expression or non-empty default-kind character literal" }
+!$omp interop init ( prefer_type( {fr("hello" }) : obj1) ! { dg-error "47: Expected '\\)'" }
+! { dg-warning "Unknown foreign runtime identifier 'hello' at \\(1\\) \\\[-Wopenmp\\\]" "" { target *-*-* } .-1 }
+
+!$omp interop init ( prefer_type( {fr(x) }) : obj1) ! { dg-error "Expected constant scalar integer expression or non-empty default-kind character literal" }
+!$omp interop init ( prefer_type( {fr(ifr_array ) }) : obj1) ! { dg-error "Expected constant scalar integer expression or non-empty default-kind character literal" }
+!$omp interop init ( prefer_type( {fr(ifr_array(1) ) }) : obj1)
 
 !$omp interop init ( prefer_type( omp_ifr_level_zero, omp_ifr_hip ) : obj1) ! OK
 !$omp interop init ( prefer_type( omp_ifr_level_zero +1 ) : obj1) ! OK
@@ -75,8 +88,10 @@  integer :: x
 !$omp interop init ( prefer_type( 4, omp_ifr_hip*4) : obj1) ! { dg-warning "Unknown foreign runtime identifier '20'" }
 !$omp interop init ( prefer_type( 4, 1, 3) : obj1)
 
-!$omp interop init ( prefer_type( {fr("cuda","sycl") }, {fr(omp_ifr_hsa,omp_ifr_level_zero)} , {attr("ompx_a") } , {fr(omp_ifr_hip) }) : obj1)
-!$omp interop init ( prefer_type( {fr("cuda","sycl"), attr("ompx_1", "ompx_2"), attr("ompx_3") }, {attr("ompx_4", "ompx_5"),fr(omp_ifr_hsa,omp_ifr_level_zero)} ) : obj1)
-!$omp interop init ( prefer_type( { fr("cuda","sycl"), attr("ompx_1") }, {fr(omp_ifr_hsa,omp_ifr_level_zero)} , {attr("ompx_a") } ) : obj1)
+!$omp interop init ( prefer_type( {fr("cuda") }, {fr(omp_ifr_hsa)} , {attr("ompx_a") } , {fr(omp_ifr_hip) }) : obj1)
+!$omp interop init ( prefer_type( {fr("cuda") }, {fr(omp_ifr_hsa,omp_ifr_level_zero)} , {attr("ompx_a") } , {fr(omp_ifr_hip) }) : obj1) ! { dg-error "65: Expected '\\)'" }
+!$omp interop init ( prefer_type( {fr("cuda",5) }, {fr(omp_ifr_hsa,omp_ifr_level_zero)} , {attr("ompx_a") } , {fr(omp_ifr_hip) }) : obj1) ! { dg-error "45: Expected '\\)' at" }
+!$omp interop init ( prefer_type( {fr("sycl"), attr("ompx_1", "ompx_2"), attr("ompx_3") }, {attr("ompx_4", "ompx_5"),fr(omp_ifr_level_zero)} ) : obj1)
+!$omp interop init ( prefer_type( { fr(5), attr("ompx_1") }, {fr(omp_ifr_hsa)} , {attr("ompx_a") } ) : obj1)
 
 end
diff --git a/gcc/testsuite/gfortran.dg/gomp/interop-2.f90 b/gcc/testsuite/gfortran.dg/gomp/interop-2.f90
index f3391bf88f0..b3130117fb2 100644
--- a/gcc/testsuite/gfortran.dg/gomp/interop-2.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/interop-2.f90
@@ -17,16 +17,50 @@  module m
  integer (omp_interop_fr_kind), parameter :: omp_ifr_hsa = 7
 end module m
 
-program main
+subroutine s(ointent)
 use m
 implicit none
+integer(omp_interop_kind), parameter :: op = 0
+integer(omp_interop_kind),intent(in) :: ointent
+integer(omp_interop_kind) :: od(5)
+integer(1) :: o1
+integer, parameter :: mykind = mod (omp_interop_kind, 100) ! remove saving the 'comes from c_int' info
+real(mykind) :: or
+
+!$omp interop init (op)      ! { dg-error "'op' at \\(1\\) in 'INIT' clause must be a scalar integer variable of 'omp_interop_kind' kind" }
+                             ! { dg-error "Object 'op' is not a variable at \\(1\\)" "" { target *-*-* } .-1 }
+!$omp interop init (ointent) ! { dg-error "'ointent' at \\(1\\) in 'INIT' clause must be definable" }
+!$omp interop init (od)      ! { dg-error "'od' at \\(1\\) in 'INIT' clause must be a scalar integer variable of 'omp_interop_kind' kind" }
+!$omp interop init (od(1))   ! { dg-error "Syntax error in OpenMP variable list" }
+!$omp interop init (o1)      ! { dg-error "'o1' at \\(1\\) in 'INIT' clause must be a scalar integer variable of 'omp_interop_kind' kind" }
+!$omp interop init (or)      ! { dg-error "'or' at \\(1\\) in 'INIT' clause must be a scalar integer variable of 'omp_interop_kind' kind" }
+
+!$omp interop use (op)      ! { dg-error "'op' at \\(1\\) in 'USE' clause must be a scalar integer variable of 'omp_interop_kind' kind" }
+                            ! { dg-error "Object 'op' is not a variable at \\(1\\)" "" { target *-*-* } .-1 }
+!$omp interop use (ointent) ! okay
+!$omp interop use (od)      ! { dg-error "'od' at \\(1\\) in 'USE' clause must be a scalar integer variable of 'omp_interop_kind' kind" }
+!$omp interop use (od(1))   ! { dg-error "Syntax error in OpenMP variable list" }
+!$omp interop use (o1)      ! { dg-error "'o1' at \\(1\\) in 'USE' clause must be a scalar integer variable of 'omp_interop_kind' kind" }
+!$omp interop use (or)      ! { dg-error "'or' at \\(1\\) in 'USE' clause must be a scalar integer variable of 'omp_interop_kind' kind" }
+
+!$omp interop destroy (op)      ! { dg-error "'op' at \\(1\\) in 'DESTROY' clause must be a scalar integer variable of 'omp_interop_kind' kind" }
+                                ! { dg-error "Object 'op' is not a variable at \\(1\\)" "" { target *-*-* } .-1 }
+!$omp interop destroy (ointent) ! { dg-error "'ointent' at \\(1\\) in 'DESTROY' clause must be definable" }
+!$omp interop destroy (od)      ! { dg-error "'od' at \\(1\\) in 'DESTROY' clause must be a scalar integer variable of 'omp_interop_kind' kind" }
+!$omp interop destroy (od(1))   ! { dg-error "Syntax error in OpenMP variable list" }
+!$omp interop destroy (o1)      ! { dg-error "'o1' at \\(1\\) in 'DESTROY' clause must be a scalar integer variable of 'omp_interop_kind' kind" }
+!$omp interop destroy (or)      ! { dg-error "'or' at \\(1\\) in 'DESTROY' clause must be a scalar integer variable of 'omp_interop_kind' kind" }
 
-!$omp requires reverse_offload
+end subroutine
+
+program main
+use m
+implicit none
 
 integer(omp_interop_kind) :: obj1, obj2, obj3, obj4, obj5
 integer :: x
 
-!$omp interop init ( prefer_type( {fr(1_"") }) : obj1) ! { dg-error "Expected scalar integer parameter or non-empty default-kind character literal" }
+!$omp interop init ( prefer_type( {fr(1_"") }) : obj1) ! { dg-error "Expected constant scalar integer expression or non-empty default-kind character literal" }
 !$omp interop init ( prefer_type( {fr(1_"hip") , attr(omp_ifr_cuda) }) : obj1) ! { dg-error "Expected default-kind character literal" }
 
 !$omp interop init ( prefer_type( {fr(1_"hip") , attr("myooption") }) : obj1) ! { dg-error "Character literal at .1. must start with 'ompx_'" }
diff --git a/gcc/testsuite/gfortran.dg/gomp/interop-3.f90 b/gcc/testsuite/gfortran.dg/gomp/interop-3.f90
index 462ed4f2e4b..02e6c903fb0 100644
--- a/gcc/testsuite/gfortran.dg/gomp/interop-3.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/interop-3.f90
@@ -21,19 +21,17 @@  program main
 use m
 implicit none
 
-!$omp requires reverse_offload
-
 integer(omp_interop_kind) :: obj1, obj2, obj3, obj4, obj5
 integer(omp_interop_kind) :: target, targetsync,prefer_type
 integer :: x
 
-!$omp interop init(obj1) init(target,targetsync,target,targetsync : obj2, obj3) nowait
+!$omp interop init(obj1) init(target,targetsync : obj2, obj3) nowait
 
-!$omp interop init(prefer_type("cu"//"da", omp_ifr_opencl, omp_ifr_level_zero, "hsa"), targetsync : obj1) &
+!$omp interop init(prefer_type(1_"cuda", omp_ifr_opencl, omp_ifr_level_zero, "hsa"), targetsync : obj1) &
 !$omp&        destroy(obj2, obj3) depend(inout: x) use(obj4, obj5) device(device_num: 0)
 
 !$omp assume contains(interop)
-  !$omp interop init(prefer_type("cu"//char(1)//"da") : obj3)  ! { dg-warning "Unknown foreign runtime identifier 'cu\\\\x01da'" }
+  !$omp interop init(prefer_type("cu da") : obj3)  ! { dg-warning "Unknown foreign runtime identifier 'cu da'" }
 !$omp end assume
 
 !$omp interop init(obj1, obj2, obj1), use(obj4) destroy(obj4)
@@ -53,7 +51,16 @@  integer :: x
 
 !$omp interop init(target, targetsync, prefer_type, obj1)
 !$omp interop init(prefer_type, obj1, target, targetsync)
-!$omp interop init(target, targetsync,target)  ! { dg-error "Symbol 'target' present on multiple clauses" }
+
+! Duplicated variable name or duplicated modifier:
+!$omp interop init(target, targetsync,target : obj1)  ! { dg-error "Duplicate 'target' at \\(1\\)" }
+!$omp interop init(target, targetsync,target)         ! { dg-error "Duplicate 'target' at \\(1\\)" }
+!$omp interop init(target : target, targetsync,target)  ! { dg-error "Symbol 'target' present on multiple clauses" }
+
+!$omp interop init(target, targetsync,targetsync : obj1)  ! { dg-error "Duplicate 'targetsync' at \\(1\\)" }
+!$omp interop init(target, targetsync,targetsync)         ! { dg-error "Duplicate 'targetsync' at \\(1\\)" }
+!$omp interop init(target : target, targetsync,targetsync)  ! { dg-error "Symbol 'targetsync' present on multiple clauses" }
+
 
 !$omp interop init(, targetsync, prefer_type, obj1, target)  ! { dg-error "Syntax error in OpenMP variable list" }
 end
diff --git a/gcc/testsuite/gfortran.dg/gomp/interop-4.f90 b/gcc/testsuite/gfortran.dg/gomp/interop-4.f90
new file mode 100644
index 00000000000..8783f4cfb5f
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/interop-4.f90
@@ -0,0 +1,56 @@ 
+! { dg-additional-options "-fdump-tree-original" }
+
+module m
+ use iso_c_binding
+ implicit none
+
+ ! The following definitions are in omp_lib, which cannot be included
+ ! in gcc/testsuite/
+ integer, parameter :: omp_interop_kind = c_intptr_t
+ integer, parameter :: omp_interop_fr_kind = c_int
+
+ integer (omp_interop_kind), parameter :: omp_interop_none = 0_omp_interop_kind
+ integer (omp_interop_fr_kind), parameter :: omp_ifr_cuda = 1
+ integer (omp_interop_fr_kind), parameter :: omp_ifr_cuda_driver = 2
+ integer (omp_interop_fr_kind), parameter :: omp_ifr_opencl = 3
+ integer (omp_interop_fr_kind), parameter :: omp_ifr_sycl = 4
+ integer (omp_interop_fr_kind), parameter :: omp_ifr_hip = 5
+ integer (omp_interop_fr_kind), parameter :: omp_ifr_level_zero = 6
+ integer (omp_interop_fr_kind), parameter :: omp_ifr_hsa = 7
+end module m
+
+subroutine s
+use m
+implicit none
+
+integer(omp_interop_kind) :: obj1, obj2, obj3, obj4, obj5, obj6, obj7
+integer :: x(6)
+
+!$omp interop init ( obj1, obj2) use (obj3) destroy(obj4) init(obj5) destroy(obj6) use(obj7) ! { dg-message "'#pragma omp interop' not yet supported" }
+! { dg-final { scan-tree-dump-times "#pragma omp interop init\\(obj1\\) init\\(obj2\\) init\\(obj5\\) use\\(obj3\\) use\\(obj7\\) destroy\\(obj4\\) destroy\\(obj6\\)\[\r\n\]" 1 "original" } }
+
+!$omp interop nowait init (targetsync : obj1, obj2) use (obj3) destroy(obj4) init(target, targetsync : obj5) destroy(obj6) use(obj7) depend(inout: x) ! { dg-message "'#pragma omp interop' not yet supported" }
+! { dg-final { scan-tree-dump-times "#pragma omp interop depend\\(inout:x\\) init\\(targetsync: obj1\\) init\\(targetsync: obj2\\) init\\(target, targetsync: obj5\\) use\\(obj3\\) use\\(obj7\\) destroy\\(obj4\\) destroy\\(obj6\\) nowait\[\r\n\]" 1 "original" } }
+
+!$omp interop init ( obj1, obj2) init (target: obj3) init(targetsync : obj4) init(target,targetsync: obj5)  ! { dg-message "'#pragma omp interop' not yet supported" }
+! { dg-final { scan-tree-dump-times "#pragma omp interop init\\(obj1\\) init\\(obj2\\) init\\(target: obj3\\) init\\(targetsync: obj4\\) init\\(target, targetsync: obj5\\)\[\r\n\]" 1 "original" } }
+
+! --------------------------------------------
+
+!$omp interop init (target, prefer_type(omp_ifr_cuda, omp_ifr_cuda+1, "hsa", "myPrivateInterop", omp_ifr_cuda-2) : obj1, obj2) init (target: obj3) init(prefer_type(omp_ifr_hip, "sycl", omp_ifr_opencl), targetsync : obj4, obj7) init(target,prefer_type("level_zero", omp_ifr_level_zero+0),targetsync: obj5)  ! { dg-message "'#pragma omp interop' not yet supported" }
+!
+! { dg-warning "Unknown foreign runtime identifier 'myPrivateInterop' at \\(1\\) \\\[-Wopenmp\\\]" "" { target *-*-* } .-2 }
+! { dg-warning "Unknown foreign runtime identifier '-1' at \\(1\\) \\\[-Wopenmp\\\]" "" { target *-*-* } .-3 }
+!
+! { dg-final { scan-tree-dump-times "#pragma omp interop init\\(prefer_type\\({fr\\(\"cuda\"\\)}, {fr\\(\"cuda_driver\"\\)}, {fr\\(\"hsa\"\\)}, {fr\\(\"<unknown>\"\\)}, {fr\\(\"<unknown>\"\\)}\\), target: obj1\\) init\\(prefer_type\\({fr\\(\"cuda\"\\)}, {fr\\(\"cuda_driver\"\\)}, {fr\\(\"hsa\"\\)}, {fr\\(\"<unknown>\"\\)}, {fr\\(\"<unknown>\"\\)}\\), target: obj2\\) init\\(target: obj3\\) init\\(prefer_type\\({fr\\(\"hip\"\\)}, {fr\\(\"sycl\"\\)}, {fr\\(\"opencl\"\\)}\\), targetsync: obj4\\) init\\(prefer_type\\({fr\\(\"hip\"\\)}, {fr\\(\"sycl\"\\)}, {fr\\(\"opencl\"\\)}\\), targetsync: obj7\\) init\\(prefer_type\\({fr\\(\"level_zero\"\\)}, {fr\\(\"level_zero\"\\)}\\), target, targetsync: obj5\\)\[\r\n\]" 1 "original" } }
+
+
+! --------------------------------------------
+
+!$omp interop init ( target, prefer_type( {fr(1_"hip"), attr("ompx_gnu_prio:1", 1_"ompx_gnu_debug")}, {attr("ompx_gnu_nicest"), attr("ompx_something")}) : obj1, obj2) init ( prefer_type( {fr("cuda")}, {fr(omp_ifr_cuda_driver), attr("ompx_nix")}, {fr("best")}), targetsync : obj3, obj4) nowait use(obj5)    ! { dg-message "'#pragma omp interop' not yet supported" }
+!
+! ! { dg-warning "Unknown foreign runtime identifier 'best' at \\(1\\) \\\[-Wopenmp\\\]" "" { target *-*-* } .-2 }
+!
+! { dg-final { scan-tree-dump-times "#pragma omp interop init\\(prefer_type\\({fr\\(\"hip\"\\),attr\\(\"ompx_gnu_prio:1\"\\),attr\\(\"ompx_gnu_debug\"\\)}, {attr\\(\"ompx_gnu_nicest\"\\),attr\\(\"ompx_something\"\\)}\\), target: obj1\\) init\\(prefer_type\\({fr\\(\"hip\"\\),attr\\(\"ompx_gnu_prio:1\"\\),attr\\(\"ompx_gnu_debug\"\\)}, {attr\\(\"ompx_gnu_nicest\"\\),attr\\(\"ompx_something\"\\)}\\), target: obj2\\) init\\(prefer_type\\({fr\\(\"cuda\"\\)}, {fr\\(\"cuda_driver\"\\),attr\\(\"ompx_nix\"\\)}, {fr\\(\"<unknown>\"\\)}\\), targetsync: obj3\\) init\\(prefer_type\\({fr\\(\"cuda\"\\)}, {fr\\(\"cuda_driver\"\\),attr\\(\"ompx_nix\"\\)}, {fr\\(\"<unknown>\"\\)}\\), targetsync: obj4\\) use\\(obj5\\) nowait\[\r\n\]" 1 "original" } }
+
+end
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index 2a91e7935ca..07d5dfa6b93 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -374,6 +374,19 @@  enum omp_clause_code {
 
   /* Range END above for: OMP_CLAUSE_SIZE */
 
+  /* OpenMP clause: destroy (variable-list ).  */
+  OMP_CLAUSE_DESTROY,
+
+  /* Range START below for: OMP_CLAUSE_INIT_PREFER_TYPE  */
+
+  /* OpenMP clause: init ( [modifier-list : ] variable-list ).  */
+  OMP_CLAUSE_INIT,
+
+  /* Range END above for: OMP_CLAUSE_INIT_PREFER_TYPE  */
+
+  /* OpenMP clause: use (variable-list ).  */
+  OMP_CLAUSE_USE,
+
   /* OpenACC clause: gang [(gang-argument-list)].
      Where
       gang-argument-list: [gang-argument-list, ] gang-argument
diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc
index eab5c60579d..e35c86812b6 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -452,6 +452,49 @@  dump_omp_iterators (pretty_printer *pp, tree iter, int spc, dump_flags_t flags)
   pp_right_paren (pp);
 }
 
+/* Dump OpenMP's prefer_type of the init clause.  */
+
+static void
+dump_omp_init_prefer_type (pretty_printer *pp, tree t)
+{
+  if (t == NULL_TREE)
+    return;
+  pp_string (pp, "prefer_type(");
+  const char *str = TREE_STRING_POINTER (t);
+  while (str[0] == (char) GOMP_INTEROP_IFR_SEPARATOR)
+    {
+      bool has_fr = false;
+      pp_character (pp, '{');
+      str++;
+      while (str[0] != (char) GOMP_INTEROP_IFR_SEPARATOR)
+	{
+	  if (has_fr)
+	    pp_character (pp, ',');
+	  has_fr = true;
+	  pp_string (pp, "fr(\"");
+	  pp_string (pp, omp_get_name_from_fr_id (str[0]));
+	  pp_string (pp, "\")");
+	  str++;
+	}
+      str++;
+      if (has_fr && str[0] != '\0')
+	pp_character (pp, ',');
+      while (str[0] != '\0')
+	{
+	  pp_string (pp, "attr(\"");
+	  pp_string (pp, str);
+	  pp_string (pp, "\")");
+	  str += strlen (str) + 1;
+	  if (str[0] != '\0')
+	    pp_character (pp, ',');
+	}
+      str++;
+      pp_character (pp, '}');
+      if (str[0] != '\0')
+	pp_string (pp, ", ");
+    }
+  pp_right_paren (pp);
+}
 
 /* Dump OMP clause CLAUSE, without following OMP_CLAUSE_CHAIN.
 
@@ -589,6 +632,44 @@  dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
       pp_right_paren (pp);
       break;
 
+    case OMP_CLAUSE_DESTROY:
+      pp_string (pp, "destroy(");
+      dump_generic_node (pp, OMP_CLAUSE_DECL (clause),
+			 spc, flags, false);
+      pp_right_paren (pp);
+      break;
+
+    case OMP_CLAUSE_INIT:
+      pp_string (pp, "init(");
+      dump_omp_init_prefer_type (pp, OMP_CLAUSE_INIT_PREFER_TYPE (clause));
+      if (OMP_CLAUSE_INIT_TARGET (clause))
+	{
+	  if (OMP_CLAUSE_INIT_PREFER_TYPE (clause))
+	    pp_string (pp, ", ");
+	  pp_string (pp, "target");
+	}
+      if (OMP_CLAUSE_INIT_TARGETSYNC (clause))
+	{
+	  if (OMP_CLAUSE_INIT_PREFER_TYPE (clause) || OMP_CLAUSE_INIT_TARGET (clause))
+	    pp_string (pp, ", ");
+	  pp_string (pp, "targetsync");
+	}
+      if (OMP_CLAUSE_INIT_PREFER_TYPE (clause)
+	  || OMP_CLAUSE_INIT_TARGET (clause)
+	  || OMP_CLAUSE_INIT_TARGETSYNC (clause))
+	pp_string (pp, ": ");
+      dump_generic_node (pp, OMP_CLAUSE_DECL (clause),
+			 spc, flags, false);
+      pp_right_paren (pp);
+      break;
+
+    case OMP_CLAUSE_USE:
+      pp_string (pp, "use(");
+      dump_generic_node (pp, OMP_CLAUSE_DECL (clause),
+			 spc, flags, false);
+      pp_right_paren (pp);
+      break;
+
     case OMP_CLAUSE_SELF:
       pp_string (pp, "self(");
       dump_generic_node (pp, OMP_CLAUSE_SELF_EXPR (clause),
@@ -3970,6 +4051,12 @@  dump_generic_node (pretty_printer *pp, tree node, int spc, dump_flags_t flags,
       dump_omp_clauses (pp, OMP_SECTIONS_CLAUSES (node), spc, flags);
       goto dump_omp_body;
 
+    case OMP_INTEROP:
+      pp_string (pp, "#pragma omp interop");
+      dump_omp_clauses (pp, OMP_INTEROP_CLAUSES (node), spc, flags);
+      is_expr = false;
+      break;
+
     case OMP_SECTION:
       pp_string (pp, "#pragma omp section");
       goto dump_omp_body;
diff --git a/gcc/tree.cc b/gcc/tree.cc
index b4c059d3b0d..51258cf5c7f 100644
--- a/gcc/tree.cc
+++ b/gcc/tree.cc
@@ -271,6 +271,9 @@  unsigned const char omp_clause_num_ops[] =
   1, /* OMP_CLAUSE_HAS_DEVICE_ADDR  */
   1, /* OMP_CLAUSE_DOACROSS  */
   2, /* OMP_CLAUSE__CACHE_  */
+  1, /* OMP_CLAUSE_DESTROY  */
+  2, /* OMP_CLAUSE_INIT  */
+  1, /* OMP_CLAUSE_USE  */
   2, /* OMP_CLAUSE_GANG  */
   1, /* OMP_CLAUSE_ASYNC  */
   1, /* OMP_CLAUSE_WAIT  */
@@ -367,6 +370,9 @@  const char * const omp_clause_code_name[] =
   "has_device_addr",
   "doacross",
   "_cache_",
+  "destroy",
+  "init",
+  "use",
   "gang",
   "async",
   "wait",
diff --git a/gcc/tree.def b/gcc/tree.def
index dd60d1ecde7..76404d598a9 100644
--- a/gcc/tree.def
+++ b/gcc/tree.def
@@ -1246,7 +1246,7 @@  DEFTREECODE (OMP_TILE, "omp_tile", tcc_statement, 7)
    Operands like for OMP_FOR.  */
 DEFTREECODE (OMP_UNROLL, "omp_unroll", tcc_statement, 7)
 
-/* OpenMP - #pragma acc loop [clause1 ... clauseN]
+/* OpenACC - #pragma acc loop [clause1 ... clauseN]
    Operands like for OMP_FOR.  */
 DEFTREECODE (OACC_LOOP, "oacc_loop", tcc_statement, 7)
 
@@ -1306,6 +1306,10 @@  DEFTREECODE (OMP_MASKED, "omp_masked", tcc_statement, 2)
    Operand 1: OMP_SCAN_CLAUSES: List of clauses.  */
 DEFTREECODE (OMP_SCAN, "omp_scan", tcc_statement, 2)
 
+/* OpenMP - #pragma omp interop [clause1 ... clauseN]
+   Operand 0: OMP_INTEROP_CLAUSES: List of clauses.  */
+DEFTREECODE (OMP_INTEROP, "omp_inteorp", tcc_statement, 1)
+
 /* OpenMP - #pragma omp section
    Operand 0: OMP_SECTION_BODY: Section body.  */
 DEFTREECODE (OMP_SECTION, "omp_section", tcc_statement, 1)
diff --git a/gcc/tree.h b/gcc/tree.h
index 7ed6d20c2e9..b47f5a2b059 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1555,6 +1555,9 @@  class auto_suppress_location_wrappers
 #define OMP_FOR_PRE_BODY(NODE)	   TREE_OPERAND (OMP_LOOPING_CHECK (NODE), 5)
 #define OMP_FOR_ORIG_DECLS(NODE)   TREE_OPERAND (OMP_LOOPING_CHECK (NODE), 6)
 
+#define OMP_INTEROP_CLAUSES(NODE)\
+  TREE_OPERAND (OMP_INTEROP_CHECK (NODE), 0)
+
 #define OMP_LOOPXFORM_CHECK(NODE) TREE_RANGE_CHECK (NODE, OMP_TILE, OMP_UNROLL)
 #define OMP_LOOPXFORM_LOWERED(NODE) \
   (OMP_LOOPXFORM_CHECK (NODE)->base.public_flag)
@@ -1824,6 +1827,15 @@  class auto_suppress_location_wrappers
 #define OMP_CLAUSE_MOTION_PRESENT(NODE) \
   (OMP_CLAUSE_RANGE_CHECK (NODE, OMP_CLAUSE_FROM, OMP_CLAUSE_TO)->base.deprecated_flag)
 
+#define OMP_CLAUSE_INIT_TARGET(NODE) \
+  (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_INIT)->base.public_flag)
+#define OMP_CLAUSE_INIT_TARGETSYNC(NODE) \
+  (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_INIT)->base.deprecated_flag)
+#define OMP_CLAUSE_INIT_PREFER_TYPE(NODE)				\
+  OMP_CLAUSE_OPERAND (OMP_CLAUSE_RANGE_CHECK (OMP_CLAUSE_CHECK (NODE),	\
+					      OMP_CLAUSE_INIT,		\
+					      OMP_CLAUSE_INIT), 1)
+
 /* Nonzero if this map clause is for array (rather than pointer) based array
    section with zero bias.  Both the non-decl OMP_CLAUSE_MAP and corresponding
    OMP_CLAUSE_MAP with GOMP_MAP_POINTER are marked with this flag.  */
diff --git a/include/gomp-constants.h b/include/gomp-constants.h
index 3091496495d..3519a8a3712 100644
--- a/include/gomp-constants.h
+++ b/include/gomp-constants.h
@@ -400,10 +400,11 @@  enum gomp_map_kind
 #define GOMP_REQUIRES_TARGET_USED           0x200
 #define GOMP_REQUIRES_SELF_MAPS             0x400
 
-/* Interop foreign-runtime data.  */
+/* Interop foreign-runtime data;
+   OpenMP defines positive values; reserve 0 and negative for GCC.  */
 #define GOMP_INTEROP_IFR_LAST	7
-#define GOMP_INTEROP_IFR_SEPARATOR -1
-#define GOMP_INTEROP_IFR_NONE -2
+#define GOMP_INTEROP_IFR_SEPARATOR -__INT8_MAX__-1
+#define GOMP_INTEROP_IFR_UNKNOWN -__INT8_MAX__
 
 /* HSA specific data structures.  */