public inbox for gcc-patches@gcc.gnu.org
 help / color / mirror / Atom feed
* [PATCH] openmp: Add support for the 'indirect' clause in C/C++
@ 2023-10-08 13:13 Kwok Cheung Yeung
  2023-10-17 13:12 ` Tobias Burnus
  0 siblings, 1 reply; 28+ messages in thread
From: Kwok Cheung Yeung @ 2023-10-08 13:13 UTC (permalink / raw)
  To: gcc-patches; +Cc: Jakub Jelinek, tobias

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

Hello

This patch adds support for the 'indirect' clause in the 'declare 
target' directive in C/C++ (Fortran to follow) and adds the necessary 
infrastructure to support indirect calls in target regions. This allows 
one to pass in pointers to functions that have been declared as indirect 
from the host to the target, then invoked via the passed-in pointer on 
the target device.

This is done by processing the functions declared as indirect in a 
similar way to regular kernels - they are added as a separate entry to 
the offload tables which are embedded into the target code by mkoffload. 
When the image is loaded, the host reads the target version of the 
offload table, then combines it with the host version to produce an 
address map. This map is then written to the device memory and a pointer 
is set to point to it.

The omp_device_lower pass now runs if any indirect functions are 
present. The pass searches for any indirect function calls, and runs a
new builtin BUILT_IN_GOMP_TARGET_MAP_INDIRECT_PTR to process the 
function pointer before making the indirect call.

The builtin (implemented by GOMP_target_map_indirect_ptr) searches 
through the address map, returning the target address if found, or the 
original address if not. I've added two search algorithms - a simple 
linear search through the map, and another which builds up a splay tree 
from the map and uses that to do the search. I've enabled the splay-tree 
version by default, but the linear search is useful for debugging 
purposes so I have kept it in.

The C++ support is currently limited to normal indirect calls - virtual 
calls on objects do not currently work. I believe the main issue is that 
the vtables are not currently copied across to the target. I have added 
some handling for OBJ_TYPE_REF to prevent the compiler from ICEing when 
it encounters a virtual call, but without the vtable this cannot work 
properly.

Tested on a x86_64 host with offloading to NVPTX and AMD GCN, and 
bootstrapped on a x86_64 host. Okay for mainline?

Thanks

Kwok

[-- Attachment #2: 0001-openmp-Add-support-for-the-indirect-clause-in-C-C.patch --]
[-- Type: text/plain, Size: 75111 bytes --]

From 46129c254990a9fff4b6d8512f04ad8fa7d61f0e Mon Sep 17 00:00:00 2001
From: Kwok Cheung Yeung <kcy@codesourcery.com>
Date: Sun, 8 Oct 2023 13:50:25 +0100
Subject: [PATCH] openmp: Add support for the 'indirect' clause in C/C++

This adds support for the 'indirect' clause in the 'declare target'
directive.  Functions declared as indirect may be called via function
pointers passed from the host in offloaded code.

Virtual calls to member functions via the object pointer in C++ are
currently not supported in target regions.

2023-10-08  Kwok Cheung Yeung  <kcy@codesourcery.com>

gcc/c-family/
	* c-attribs.cc (c_common_attribute_table): Add attribute for
	indirect functions.
	* c-pragma.h (enum parma_omp_clause): Add entry for indirect clause.

gcc/c/
	* c-decl.cc (c_decl_attributes): Add attribute for indirect
	functions.
	* c-lang.h (c_omp_declare_target_attr): Add indirect field.
	* c-parser.cc (c_parser_omp_clause_name): Handle indirect clause.
	(c_parser_omp_clause_indirect): New.
	(c_parser_omp_all_clauses): Handle indirect clause.
	(OMP_DECLARE_TARGET_CLAUSE_MASK): Add indirect clause to mask.
	(c_parser_omp_declare_target): Handle indirect clause.
	(OMP_BEGIN_DECLARE_TARGET_CLAUSE_MASK): Add indirect clause to mask.
	(c_parser_omp_begin): Handle indirect clause.
	* c-typeck.cc (c_finish_omp_clauses): Handle indirect clause.

gcc/cp/
	* cp-tree.h (cp_omp_declare_target_attr): Add indirect field.
	* decl2.cc (cplus_decl_attributes): Add attribute for indirect
	functions.
	* parser.cc (cp_parser_omp_clause_name): Handle indirect clause.
	(cp_parser_omp_clause_indirect): New.
	(cp_parser_omp_all_clauses): Handle indirect clause.
	(handle_omp_declare_target_clause): Add extra parameter.  Add
	indirect attribute for indirect functions.
	(OMP_DECLARE_TARGET_CLAUSE_MASK): Add indirect clause to mask.
	(cp_parser_omp_declare_target): Handle indirect clause.
	(OMP_BEGIN_DECLARE_TARGET_CLAUSE_MASK): Add indirect clause to mask.
	(cp_parser_omp_begin): Handle indirect clause.
	* semantics.cc (finish_omp_clauses): Handle indirect clause.

gcc/
	* lto-cgraph.cc (enum LTO_symtab_tags): Add tag for indirect
	functions.
	(output_offload_tables): Write indirect functions.
	(input_offload_tables): read indirect functions.
	* lto-section-names.h (OFFLOAD_IND_FUNC_TABLE_SECTION_NAME): New.
	* omp-builtins.def (BUILT_IN_GOMP_TARGET_MAP_INDIRECT_PTR): New.
	* omp-offload.cc (offload_ind_funcs): New.
	(omp_discover_implicit_declare_target): Add functions marked with
	'omp declare target indirect' to indirect functions list.
	(omp_finish_file): Add indirect functions to section for offload
	indirect functions.
	(execute_omp_device_lower): Redirect indirect calls on target by
	passing function pointer to BUILT_IN_GOMP_TARGET_MAP_INDIRECT_PTR.
	(pass_omp_device_lower::gate): Run pass_omp_device_lower if
	indirect functions are present.
	* omp-offload.h (offload_ind_funcs): New.
	* tree-core.h (omp_clause_code): Add OMP_CLAUSE_INDIRECT.
	* tree.cc (omp_clause_num_ops): Add entry for OMP_CLAUSE_INDIRECT.
	(omp_clause_code_name): Likewise.
	* tree.h (OMP_CLAUSE_INDIRECT_EXPR): New.
	* config/gcn/mkoffload.cc (process_asm): Process offload_ind_funcs
	section.  Count number of indirect functions.
	(process_obj): Emit number of indirect functions.
	* config/nvptx/mkoffload.cc (ind_func_ids, ind_funcs_tail): New.
	(process): Emit offload_ind_func_table in PTX code.  Emit indirect
	function names and count in image.
	* config/nvptx/nvptx.cc (nvptx_record_offload_symbol): Mark
	indirect functions in PTX code with IND_FUNC_MAP.

gcc/testsuite/
	* c-c++-common/gomp/declare-target-indirect-1.c: New.
	* c-c++-common/gomp/declare-target-indirect-2.c: New.

libgcc/
	* offloadstuff.c (OFFLOAD_IND_FUNC_TABLE_SECTION_NAME): New.
	(__offload_ind_func_table): New.
	(__offload_ind_funcs_end): New.
	(__OFFLOAD_TABLE__): Add entries for indirect functions.

libgomp/
	* Makefile.am (libgomp_la_SOURCES): Add target-indirect.c.
	* Makefile.in: Regenerate.
	* libgomp-plugin.h (GOMP_INDIRECT_ADDR_MAP): New define.
	(GOMP_OFFLOAD_load_image): Add extra argument.
	* libgomp.h (struct indirect_splay_tree_key_s): New.
	(indirect_splay_tree_node, indirect_splay_tree,
	indirect_splay_tree_key): New.
	(indirect_splay_compare): New.
	* libgomp.map (GOMP_5.1.1): Add GOMP_target_map_indirect_ptr.
	* libgomp.texi (OpenMP 5.1): Update documentation on indirect
	calls in target region and on indirect clause.
	* libgomp_g.h (GOMP_target_map_indirect_ptr): Add prototype.
	* oacc-host.c (host_load_image): Add extra argument.
	* target.c (gomp_load_image_to_device): Pass host indirect
	functions table to load_image_func.
	* config/accel/target-indirect.c: New.
	* config/linux/target-indirect.c: New.
	* config/gcn/team.c (build_indirect_map): Add prototype.
	(gomp_gcn_enter_kernel): Initialize support for indirect
	function calls on GCN target.
	* config/nvptx/team.c (build_indirect_map): Add prototype.
	(gomp_nvptx_main): Initialize support for indirect function
	calls on NVPTX target.
	* plugin/plugin-gcn.c (struct gcn_image_desc): Add field for
	indirect functions count.
	(GOMP_OFFLOAD_load_image): Add extra argument.  Build address
	translation table and copy it to target memory.
	* plugin/plugin-nvptx.c (nvptx_tdata): Add field for indirect
	functions count.
	(GOMP_OFFLOAD_load_image): Add extra argument.  Build address
	translation table and copy it to target memory.
	* testsuite/libgomp.c-c++-common/declare-target-indirect-1.c: New.
	* testsuite/libgomp.c-c++-common/declare-target-indirect-2.c: New.
---
 gcc/c-family/c-attribs.cc                     |   2 +
 gcc/c-family/c-pragma.h                       |   1 +
 gcc/c/c-decl.cc                               |   8 ++
 gcc/c/c-lang.h                                |   1 +
 gcc/c/c-parser.cc                             |  89 +++++++++++--
 gcc/c/c-typeck.cc                             |   1 +
 gcc/config/gcn/mkoffload.cc                   |  29 +++-
 gcc/config/nvptx/mkoffload.cc                 |  87 +++++++++++-
 gcc/config/nvptx/nvptx.cc                     |   6 +-
 gcc/cp/cp-tree.h                              |   1 +
 gcc/cp/decl2.cc                               |   6 +
 gcc/cp/parser.cc                              | 100 ++++++++++++--
 gcc/cp/semantics.cc                           |   1 +
 gcc/lto-cgraph.cc                             |  27 ++++
 gcc/lto-section-names.h                       |   1 +
 gcc/omp-builtins.def                          |   3 +
 gcc/omp-offload.cc                            |  77 +++++++++--
 gcc/omp-offload.h                             |   1 +
 .../gomp/declare-target-indirect-1.c          |  51 +++++++
 .../gomp/declare-target-indirect-2.c          |  20 +++
 gcc/tree-core.h                               |   3 +
 gcc/tree.cc                                   |   2 +
 gcc/tree.h                                    |   4 +
 libgcc/offloadstuff.c                         |  12 +-
 libgomp/Makefile.am                           |   2 +-
 libgomp/Makefile.in                           |   5 +-
 libgomp/config/accel/target-indirect.c        | 126 ++++++++++++++++++
 libgomp/config/gcn/team.c                     |   4 +
 libgomp/config/linux/target-indirect.c        |  31 +++++
 libgomp/config/nvptx/team.c                   |   5 +
 libgomp/libgomp-plugin.h                      |   5 +-
 libgomp/libgomp.h                             |  23 ++++
 libgomp/libgomp.map                           |   1 +
 libgomp/libgomp.texi                          |   4 +-
 libgomp/libgomp_g.h                           |   1 +
 libgomp/oacc-host.c                           |   3 +-
 libgomp/plugin/plugin-gcn.c                   |  87 +++++++++++-
 libgomp/plugin/plugin-nvptx.c                 |  62 ++++++++-
 libgomp/target.c                              |  11 +-
 .../declare-target-indirect-1.c               |  21 +++
 .../declare-target-indirect-2.c               |  33 +++++
 41 files changed, 909 insertions(+), 48 deletions(-)
 create mode 100644 gcc/testsuite/c-c++-common/gomp/declare-target-indirect-1.c
 create mode 100644 gcc/testsuite/c-c++-common/gomp/declare-target-indirect-2.c
 create mode 100644 libgomp/config/accel/target-indirect.c
 create mode 100644 libgomp/config/linux/target-indirect.c
 create mode 100644 libgomp/testsuite/libgomp.c-c++-common/declare-target-indirect-1.c
 create mode 100644 libgomp/testsuite/libgomp.c-c++-common/declare-target-indirect-2.c

diff --git a/gcc/c-family/c-attribs.cc b/gcc/c-family/c-attribs.cc
index dca7548b2c6..cb409216f69 100644
--- a/gcc/c-family/c-attribs.cc
+++ b/gcc/c-family/c-attribs.cc
@@ -518,6 +518,8 @@ const struct attribute_spec c_common_attribute_table[] =
 			      handle_omp_declare_target_attribute, NULL },
   { "omp declare target implicit", 0, 0, true, false, false, false,
 			      handle_omp_declare_target_attribute, NULL },
+  { "omp declare target indirect", 0, 0, true, false, false, false,
+			      handle_omp_declare_target_attribute, NULL },
   { "omp declare target host", 0, 0, true, false, false, false,
 			      handle_omp_declare_target_attribute, NULL },
   { "omp declare target nohost", 0, 0, true, false, false, false,
diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index 603c5151978..902a924360b 100644
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -125,6 +125,7 @@ enum pragma_omp_clause {
   PRAGMA_OMP_CLAUSE_IF,
   PRAGMA_OMP_CLAUSE_IN_REDUCTION,
   PRAGMA_OMP_CLAUSE_INBRANCH,
+  PRAGMA_OMP_CLAUSE_INDIRECT,
   PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR,
   PRAGMA_OMP_CLAUSE_LASTPRIVATE,
   PRAGMA_OMP_CLAUSE_LINEAR,
diff --git a/gcc/c/c-decl.cc b/gcc/c/c-decl.cc
index 5822faf01b4..f0bc1c65621 100644
--- a/gcc/c/c-decl.cc
+++ b/gcc/c/c-decl.cc
@@ -5331,6 +5331,14 @@ c_decl_attributes (tree *node, tree attributes, int flags)
 	    attributes
 	      = tree_cons (get_identifier ("omp declare target nohost"),
 			   NULL_TREE, attributes);
+
+	  int indirect
+	    = current_omp_declare_target_attribute->last ().indirect;
+	  if (indirect && !lookup_attribute ("omp declare target indirect",
+					     attributes))
+	    attributes
+	      = tree_cons (get_identifier ("omp declare target indirect"),
+			   NULL_TREE, attributes);
 	}
     }
 
diff --git a/gcc/c/c-lang.h b/gcc/c/c-lang.h
index 4fea11855f1..cb13e34e80e 100644
--- a/gcc/c/c-lang.h
+++ b/gcc/c/c-lang.h
@@ -62,6 +62,7 @@ struct GTY(()) language_function {
 
 struct GTY(()) c_omp_declare_target_attr {
   int device_type;
+  int indirect;
 };
 
 /* If non-zero, implicit "omp declare target" attribute is added into the
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 0d468b86bd8..50120a10a5f 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -13975,6 +13975,8 @@ c_parser_omp_clause_name (c_parser *parser)
 	    result = PRAGMA_OMP_CLAUSE_IN_REDUCTION;
 	  else if (!strcmp ("inbranch", p))
 	    result = PRAGMA_OMP_CLAUSE_INBRANCH;
+	  else if (!strcmp ("indirect", p))
+	    result = PRAGMA_OMP_CLAUSE_INDIRECT;
 	  else if (!strcmp ("independent", p))
 	    result = PRAGMA_OACC_CLAUSE_INDEPENDENT;
 	  else if (!strcmp ("is_device_ptr", p))
@@ -14835,6 +14837,47 @@ c_parser_omp_clause_final (c_parser *parser, tree list)
   return list;
 }
 
+/* OpenMP 5.1:
+   indirect [( expression )]
+*/
+
+static tree
+c_parser_omp_clause_indirect (c_parser *parser, tree list)
+{
+  location_t location = c_parser_peek_token (parser)->location;
+  tree t;
+
+  if (c_parser_peek_token (parser)->type == CPP_OPEN_PAREN)
+    {
+      matching_parens parens;
+      if (!parens.require_open (parser))
+	return list;
+
+      location_t loc = c_parser_peek_token (parser)->location;
+      c_expr expr = c_parser_expr_no_commas (parser, NULL);
+      expr = convert_lvalue_to_rvalue (loc, expr, true, true);
+      t = c_objc_common_truthvalue_conversion (loc, expr.value);
+      t = c_fully_fold (t, false, NULL);
+      if (!INTEGRAL_TYPE_P (TREE_TYPE (t))
+	  || TREE_CODE (t) != INTEGER_CST)
+	{
+	  c_parser_error (parser, "expected constant integer expression");
+	  return list;
+	}
+      parens.skip_until_found_close (parser);
+    }
+  else
+    t = integer_one_node;
+
+  check_no_duplicate_clause (list, OMP_CLAUSE_INDIRECT, "indirect");
+
+  tree c = build_omp_clause (location, OMP_CLAUSE_INDIRECT);
+  OMP_CLAUSE_INDIRECT_EXPR (c) = t;
+  OMP_CLAUSE_CHAIN (c) = list;
+
+  return c;
+}
+
 /* OpenACC, OpenMP 2.5:
    if ( expression )
 
@@ -18352,6 +18395,10 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask,
 					     true, clauses);
 	  c_name = "in_reduction";
 	  break;
+	case PRAGMA_OMP_CLAUSE_INDIRECT:
+	  clauses = c_parser_omp_clause_indirect (parser, clauses);
+	  c_name = "indirect";
+	  break;
 	case PRAGMA_OMP_CLAUSE_LASTPRIVATE:
 	  clauses = c_parser_omp_clause_lastprivate (parser, clauses);
 	  c_name = "lastprivate";
@@ -23748,13 +23795,15 @@ c_finish_omp_declare_simd (c_parser *parser, tree fndecl, tree parms,
 	( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_TO)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_ENTER)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_LINK)		\
-	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE_TYPE))
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE_TYPE)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_INDIRECT))
 
 static void
 c_parser_omp_declare_target (c_parser *parser)
 {
   tree clauses = NULL_TREE;
   int device_type = 0;
+  bool indirect = false;
   bool only_device_type = true;
   if (c_parser_next_token_is (parser, CPP_NAME)
       || (c_parser_next_token_is (parser, CPP_COMMA)
@@ -23771,16 +23820,21 @@ c_parser_omp_declare_target (c_parser *parser)
   else
     {
       c_parser_skip_to_pragma_eol (parser);
-      c_omp_declare_target_attr attr = { -1 };
+      c_omp_declare_target_attr attr = { -1, 0 };
       vec_safe_push (current_omp_declare_target_attribute, attr);
       return;
     }
-  for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
-    if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEVICE_TYPE)
-      device_type |= OMP_CLAUSE_DEVICE_TYPE_KIND (c);
   for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
     {
       if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEVICE_TYPE)
+	device_type |= OMP_CLAUSE_DEVICE_TYPE_KIND (c);
+      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_INDIRECT)
+	indirect |= !integer_zerop (OMP_CLAUSE_INDIRECT_EXPR (c));
+    }
+  for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+    {
+      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEVICE_TYPE
+	  || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_INDIRECT)
 	continue;
       tree t = OMP_CLAUSE_DECL (c), id;
       tree at1 = lookup_attribute ("omp declare target", DECL_ATTRIBUTES (t));
@@ -23848,6 +23902,17 @@ c_parser_omp_declare_target (c_parser *parser)
 		= tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (t));
 	    }
 	}
+      if (indirect)
+	{
+	  tree at4 = lookup_attribute ("omp declare target indirect",
+				       DECL_ATTRIBUTES (t));
+	  if (at4 == NULL_TREE)
+	    {
+	      id = get_identifier ("omp declare target indirect");
+	      DECL_ATTRIBUTES (t)
+		= tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (t));
+	    }
+	}
     }
   if (device_type && only_device_type)
     error_at (OMP_CLAUSE_LOCATION (clauses),
@@ -23860,7 +23925,8 @@ c_parser_omp_declare_target (c_parser *parser)
    #pragma omp begin declare target clauses[optseq] new-line  */
 
 #define OMP_BEGIN_DECLARE_TARGET_CLAUSE_MASK			\
-	(OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE_TYPE)
+	( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE_TYPE)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_INDIRECT))
 
 static void
 c_parser_omp_begin (c_parser *parser)
@@ -23883,10 +23949,15 @@ c_parser_omp_begin (c_parser *parser)
 					OMP_BEGIN_DECLARE_TARGET_CLAUSE_MASK,
 					"#pragma omp begin declare target");
 	  int device_type = 0;
+	  int indirect = 0;
 	  for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
-	    if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEVICE_TYPE)
-	      device_type |= OMP_CLAUSE_DEVICE_TYPE_KIND (c);
-	  c_omp_declare_target_attr attr = { device_type };
+	    {
+	      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEVICE_TYPE)
+		device_type |= OMP_CLAUSE_DEVICE_TYPE_KIND (c);
+	      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_INDIRECT)
+		indirect |= !integer_zerop (OMP_CLAUSE_INDIRECT_EXPR (c));
+	    }
+	  c_omp_declare_target_attr attr = { device_type, indirect };
 	  vec_safe_push (current_omp_declare_target_attribute, attr);
 	}
       else
diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc
index e55e887da14..ed889dd353a 100644
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -15880,6 +15880,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	case OMP_CLAUSE_IF_PRESENT:
 	case OMP_CLAUSE_FINALIZE:
 	case OMP_CLAUSE_NOHOST:
+	case OMP_CLAUSE_INDIRECT:
 	  pc = &OMP_CLAUSE_CHAIN (c);
 	  continue;
 
diff --git a/gcc/config/gcn/mkoffload.cc b/gcc/config/gcn/mkoffload.cc
index 8b608bf024e..6a109c3a926 100644
--- a/gcc/config/gcn/mkoffload.cc
+++ b/gcc/config/gcn/mkoffload.cc
@@ -477,7 +477,8 @@ copy_early_debug_info (const char *infile, const char *outfile)
 static void
 process_asm (FILE *in, FILE *out, FILE *cfile)
 {
-  int fn_count = 0, var_count = 0, dims_count = 0, regcount_count = 0;
+  int fn_count = 0, var_count = 0, ind_fn_count = 0;
+  int dims_count = 0, regcount_count = 0;
   struct obstack fns_os, dims_os, regcounts_os;
   obstack_init (&fns_os);
   obstack_init (&dims_os);
@@ -506,7 +507,8 @@ process_asm (FILE *in, FILE *out, FILE *cfile)
     { IN_CODE,
       IN_METADATA,
       IN_VARS,
-      IN_FUNCS
+      IN_FUNCS,
+      IN_IND_FUNCS,
     } state = IN_CODE;
   while (fgets (buf, sizeof (buf), in))
     {
@@ -568,6 +570,17 @@ process_asm (FILE *in, FILE *out, FILE *cfile)
 	      }
 	    break;
 	  }
+	case IN_IND_FUNCS:
+	  {
+	    char *funcname;
+	    if (sscanf (buf, "\t.8byte\t%ms\n", &funcname))
+	      {
+		fputs (buf, out);
+		ind_fn_count++;
+		continue;
+	      }
+	    break;
+	  }
 	}
 
       char dummy;
@@ -595,6 +608,15 @@ process_asm (FILE *in, FILE *out, FILE *cfile)
 		 ".offload_func_table:\n",
 		 out);
 	}
+      else if (sscanf (buf, " .section .gnu.offload_ind_funcs%c", &dummy) > 0)
+	{
+	  state = IN_IND_FUNCS;
+	  fputs (buf, out);
+	  fputs ("\t.global .offload_ind_func_table\n"
+		 "\t.type .offload_ind_func_table, @object\n"
+		 ".offload_ind_func_table:\n",
+		 out);
+	}
       else if (sscanf (buf, " .amdgpu_metadata%c", &dummy) > 0)
 	{
 	  state = IN_METADATA;
@@ -632,6 +654,7 @@ process_asm (FILE *in, FILE *out, FILE *cfile)
   fprintf (cfile, "#include <stdbool.h>\n\n");
 
   fprintf (cfile, "static const int gcn_num_vars = %d;\n\n", var_count);
+  fprintf (cfile, "static const int gcn_num_ind_funcs = %d;\n\n", ind_fn_count);
 
   /* Dump out function idents.  */
   fprintf (cfile, "static const struct hsa_kernel_description {\n"
@@ -726,12 +749,14 @@ process_obj (FILE *in, FILE *cfile, uint32_t omp_requires)
 	   "  const struct gcn_image *gcn_image;\n"
 	   "  unsigned kernel_count;\n"
 	   "  const struct hsa_kernel_description *kernel_infos;\n"
+	   "  unsigned ind_func_count;\n"
 	   "  unsigned global_variable_count;\n"
 	   "} gcn_data = {\n"
 	   "  %d,\n"
 	   "  &gcn_image,\n"
 	   "  sizeof (gcn_kernels) / sizeof (gcn_kernels[0]),\n"
 	   "  gcn_kernels,\n"
+	   "  gcn_num_ind_funcs,\n"
 	   "  gcn_num_vars\n"
 	   "};\n\n", omp_requires);
 
diff --git a/gcc/config/nvptx/mkoffload.cc b/gcc/config/nvptx/mkoffload.cc
index aaea9fb320d..fb75ca090df 100644
--- a/gcc/config/nvptx/mkoffload.cc
+++ b/gcc/config/nvptx/mkoffload.cc
@@ -51,6 +51,7 @@ struct id_map
 };
 
 static id_map *func_ids, **funcs_tail = &func_ids;
+static id_map *ind_func_ids, **ind_funcs_tail = &ind_func_ids;
 static id_map *var_ids, **vars_tail = &var_ids;
 
 /* Files to unlink.  */
@@ -302,6 +303,11 @@ process (FILE *in, FILE *out, uint32_t omp_requires)
 		      output_fn_ptr = true;
 		      record_id (input + i + 9, &funcs_tail);
 		    }
+		  else if (startswith (input + i, "IND_FUNC_MAP "))
+		    {
+		      output_fn_ptr = true;
+		      record_id (input + i + 13, &ind_funcs_tail);
+		    }
 		  else
 		    abort ();
 		  /* Skip to next line. */
@@ -422,6 +428,77 @@ process (FILE *in, FILE *out, uint32_t omp_requires)
       fprintf (out, "};\\n\";\n\n");
     }
 
+  if (ind_func_ids)
+    {
+      const char needle[] = "// BEGIN GLOBAL FUNCTION DECL: ";
+
+      fprintf (out, "static const char ptx_code_%u[] =\n", obj_count++);
+      fprintf (out, "\t\".version ");
+      for (size_t i = 0; version[i] != '\0' && version[i] != '\n'; i++)
+	fputc (version[i], out);
+      fprintf (out, "\"\n\t\".target sm_");
+      for (size_t i = 0; sm_ver[i] != '\0' && sm_ver[i] != '\n'; i++)
+	fputc (sm_ver[i], out);
+      fprintf (out, "\"\n\t\".file 2 \\\"<dummy>\\\"\"\n");
+
+      /* WORKAROUND - see PR 108098
+	 It seems as if older CUDA JIT compiler optimizes the function pointers
+	 in offload_func_table to NULL, which can be prevented by adding a
+	 dummy procedure. With CUDA 11.1, it seems to work fine without
+	 workaround while CUDA 10.2 as some ancient version have need the
+	 workaround. Assuming CUDA 11.0 fixes it, emitting it could be
+	 restricted to 'if (sm_ver2[0] < 8 && version2[0] < 7)' as sm_80 and
+	 PTX ISA 7.0 are new in CUDA 11.0; for 11.1 it would be sm_86 and
+	 PTX ISA 7.1.  */
+      fprintf (out, "\n\t\".func __dummy$func2 ( );\"\n");
+      fprintf (out, "\t\".func __dummy$func2 ( )\"\n");
+      fprintf (out, "\t\"{\"\n");
+      fprintf (out, "\t\"}\"\n");
+
+      size_t fidx = 0;
+      for (id = ind_func_ids; id; id = id->next)
+	{
+	  fprintf (out, "\t\".extern ");
+	  const char *p = input + file_idx[fidx];
+	  while (true)
+	    {
+	      p = strstr (p, needle);
+	      if (!p)
+		{
+		  fidx++;
+		  if (fidx >= file_cnt)
+		    break;
+		  p = input + file_idx[fidx];
+		  continue;
+		}
+	      p += strlen (needle);
+	      if (!startswith (p, id->ptx_name))
+		continue;
+	      p += strlen (id->ptx_name);
+	      if (*p != '\n')
+		continue;
+	      p++;
+	      /* Skip over any directives.  */
+	      while (!startswith (p, ".func"))
+		while (*p++ != ' ');
+	      for (; *p != '\0' && *p != '\n'; p++)
+		fputc (*p, out);
+	      break;
+	    }
+	  fprintf (out, "\"\n");
+	  if (fidx == file_cnt)
+	    fatal_error (input_location,
+			 "Cannot find function declaration for %qs",
+			 id->ptx_name);
+	}
+
+      fprintf (out, "\t\".visible .global .align 8 .u64 "
+		    "$offload_ind_func_table[] = {");
+      for (comma = "", id = ind_func_ids; id; comma = ",", id = id->next)
+	fprintf (out, "%s\"\n\t\t\"%s", comma, id->ptx_name);
+      fprintf (out, "};\\n\";\n\n");
+    }
+
   /* Dump out array of pointers to ptx object strings.  */
   fprintf (out, "static const struct ptx_obj {\n"
 	   "  const char *code;\n"
@@ -447,6 +524,12 @@ process (FILE *in, FILE *out, uint32_t omp_requires)
 	     id->dim ? id->dim : "");
   fprintf (out, "\n};\n\n");
 
+  /* Dump out indirect function idents.  */
+  fprintf (out, "static const char *const ind_func_mappings[] = {");
+  for (comma = "", id = ind_func_ids; id; comma = ",", id = id->next)
+    fprintf (out, "%s\n\t\"%s\"", comma, id->ptx_name);
+  fprintf (out, "\n};\n\n");
+
   fprintf (out,
 	   "static const struct nvptx_data {\n"
 	   "  uintptr_t omp_requires_mask;\n"
@@ -456,12 +539,14 @@ process (FILE *in, FILE *out, uint32_t omp_requires)
 	   "  unsigned var_num;\n"
 	   "  const struct nvptx_fn *fn_names;\n"
 	   "  unsigned fn_num;\n"
+	   "  unsigned ind_fn_num;\n"
 	   "} nvptx_data = {\n"
 	   "  %d, ptx_objs, sizeof (ptx_objs) / sizeof (ptx_objs[0]),\n"
 	   "  var_mappings,"
 	   "  sizeof (var_mappings) / sizeof (var_mappings[0]),\n"
 	   "  func_mappings,"
-	   "  sizeof (func_mappings) / sizeof (func_mappings[0])\n"
+	   "  sizeof (func_mappings) / sizeof (func_mappings[0]),\n"
+	   "  sizeof (ind_func_mappings) / sizeof (ind_func_mappings[0])\n"
 	   "};\n\n", omp_requires);
 
   fprintf (out, "#ifdef __cplusplus\n"
diff --git a/gcc/config/nvptx/nvptx.cc b/gcc/config/nvptx/nvptx.cc
index edef39fb5e1..eaff812a520 100644
--- a/gcc/config/nvptx/nvptx.cc
+++ b/gcc/config/nvptx/nvptx.cc
@@ -5918,7 +5918,11 @@ nvptx_record_offload_symbol (tree decl)
 	/* OpenMP offloading does not set this attribute.  */
 	tree dims = attr ? TREE_VALUE (attr) : NULL_TREE;
 
-	fprintf (asm_out_file, "//:FUNC_MAP \"%s\"",
+	fprintf (asm_out_file, "//:");
+	if (lookup_attribute ("omp declare target indirect",
+			      DECL_ATTRIBUTES (decl)))
+	  fprintf (asm_out_file, "IND_");
+	fprintf (asm_out_file, "FUNC_MAP \"%s\"",
 		 IDENTIFIER_POINTER (DECL_ASSEMBLER_NAME (decl)));
 
 	for (; dims; dims = TREE_CHAIN (dims))
diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h
index 6e34952da99..8d0ad38c1ab 100644
--- a/gcc/cp/cp-tree.h
+++ b/gcc/cp/cp-tree.h
@@ -1831,6 +1831,7 @@ union GTY((desc ("cp_tree_node_structure (&%h)"),
 struct GTY(()) cp_omp_declare_target_attr {
   bool attr_syntax;
   int device_type;
+  bool indirect;
 };
 
 struct GTY(()) cp_omp_begin_assumes_data {
diff --git a/gcc/cp/decl2.cc b/gcc/cp/decl2.cc
index 344e19ec98b..f89d32950e9 100644
--- a/gcc/cp/decl2.cc
+++ b/gcc/cp/decl2.cc
@@ -1768,6 +1768,12 @@ cplus_decl_attributes (tree *decl, tree attributes, int flags)
 		attributes
 		  = tree_cons (get_identifier ("omp declare target nohost"),
 			       NULL_TREE, attributes);
+	      if (last.indirect
+		  && !lookup_attribute ("omp declare target indirect",
+					attributes))
+		attributes
+		  = tree_cons (get_identifier ("omp declare target indirect"),
+			       NULL_TREE, attributes);
 	    }
 	}
     }
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index f3abae716fe..8ca155a29f8 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -37516,6 +37516,8 @@ cp_parser_omp_clause_name (cp_parser *parser)
 	    result = PRAGMA_OMP_CLAUSE_IN_REDUCTION;
 	  else if (!strcmp ("inbranch", p))
 	    result = PRAGMA_OMP_CLAUSE_INBRANCH;
+	  else if (!strcmp ("indirect", p))
+	    result = PRAGMA_OMP_CLAUSE_INDIRECT;
 	  else if (!strcmp ("independent", p))
 	    result = PRAGMA_OACC_CLAUSE_INDEPENDENT;
 	  else if (!strcmp ("is_device_ptr", p))
@@ -38548,6 +38550,50 @@ cp_parser_omp_clause_final (cp_parser *parser, tree list, location_t location)
   return c;
 }
 
+/* OpenMP 5.1:
+   indirect [( expression )]
+*/
+
+static tree
+cp_parser_omp_clause_indirect (cp_parser *parser, tree list,
+			       location_t location)
+{
+  tree t;
+
+  if (cp_lexer_next_token_is (parser->lexer, CPP_OPEN_PAREN))
+    {
+      matching_parens parens;
+      if (!parens.require_open (parser))
+	return list;
+
+      t = cp_parser_assignment_expression (parser);
+
+      if (t != error_mark_node)
+	{
+	  t = fold_non_dependent_expr (t);
+	  if (!value_dependent_expression_p (t)
+	      && (!INTEGRAL_TYPE_P (TREE_TYPE (t))
+		  || !tree_fits_shwi_p (t)))
+	    error_at (location, "expected constant integer expression");
+	}
+      if (t == error_mark_node
+	  || !parens.require_close (parser))
+	cp_parser_skip_to_closing_parenthesis (parser, /*recovering=*/true,
+					       /*or_comma=*/false,
+					       /*consume_paren=*/true);
+    }
+  else
+    t = integer_one_node;
+
+  check_no_duplicate_clause (list, OMP_CLAUSE_INDIRECT, "indirect", location);
+
+  tree c = build_omp_clause (location, OMP_CLAUSE_INDIRECT);
+  OMP_CLAUSE_INDIRECT_EXPR (c) = t;
+  OMP_CLAUSE_CHAIN (c) = list;
+
+  return c;
+}
+
 /* OpenMP 2.5:
    if ( expression )
 
@@ -41572,6 +41618,11 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask,
 					      true, clauses);
 	  c_name = "in_reduction";
 	  break;
+	case PRAGMA_OMP_CLAUSE_INDIRECT:
+	  clauses = cp_parser_omp_clause_indirect (parser, clauses,
+						   token->location);
+	  c_name = "indirect";
+	  break;
 	case PRAGMA_OMP_CLAUSE_LASTPRIVATE:
 	  clauses = cp_parser_omp_clause_lastprivate (parser, clauses);
 	  c_name = "lastprivate";
@@ -48109,7 +48160,8 @@ cp_maybe_parse_omp_decl (tree decl, tree d)
    on #pragma omp declare target.  Return false if errors were reported.  */
 
 static bool
-handle_omp_declare_target_clause (tree c, tree t, int device_type)
+handle_omp_declare_target_clause (tree c, tree t, int device_type,
+				  bool indirect)
 {
   tree at1 = lookup_attribute ("omp declare target", DECL_ATTRIBUTES (t));
   tree at2 = lookup_attribute ("omp declare target link", DECL_ATTRIBUTES (t));
@@ -48173,6 +48225,17 @@ handle_omp_declare_target_clause (tree c, tree t, int device_type)
 	  DECL_ATTRIBUTES (t) = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (t));
 	}
     }
+  if (indirect)
+    {
+      tree at4 = lookup_attribute ("omp declare target indirect",
+				   DECL_ATTRIBUTES (t));
+      if (at4 == NULL_TREE)
+	{
+	  id = get_identifier ("omp declare target indirect");
+	  DECL_ATTRIBUTES (t)
+	    = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (t));
+	}
+    }
   return true;
 }
 
@@ -48190,13 +48253,15 @@ handle_omp_declare_target_clause (tree c, tree t, int device_type)
 	( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_TO)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_ENTER)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_LINK)		\
-	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE_TYPE))
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE_TYPE)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_INDIRECT))
 
 static void
 cp_parser_omp_declare_target (cp_parser *parser, cp_token *pragma_tok)
 {
   tree clauses = NULL_TREE;
   int device_type = 0;
+  bool indirect = false;
   bool only_device_type = true;
   if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)
       || (cp_lexer_next_token_is (parser->lexer, CPP_COMMA)
@@ -48215,21 +48280,26 @@ cp_parser_omp_declare_target (cp_parser *parser, cp_token *pragma_tok)
   else
     {
       cp_omp_declare_target_attr a
-	= { parser->lexer->in_omp_attribute_pragma, -1 };
+	= { parser->lexer->in_omp_attribute_pragma, -1, false };
       vec_safe_push (scope_chain->omp_declare_target_attribute, a);
       cp_parser_require_pragma_eol (parser, pragma_tok);
       return;
     }
-  for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
-    if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEVICE_TYPE)
-      device_type |= OMP_CLAUSE_DEVICE_TYPE_KIND (c);
   for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
     {
       if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEVICE_TYPE)
+	device_type |= OMP_CLAUSE_DEVICE_TYPE_KIND (c);
+      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_INDIRECT)
+	indirect |= !integer_zerop (OMP_CLAUSE_INDIRECT_EXPR (c));
+    }
+  for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+    {
+      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEVICE_TYPE
+	  || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_INDIRECT)
 	continue;
       tree t = OMP_CLAUSE_DECL (c);
       only_device_type = false;
-      if (!handle_omp_declare_target_clause (c, t, device_type))
+      if (!handle_omp_declare_target_clause (c, t, device_type, indirect))
 	continue;
       if (VAR_OR_FUNCTION_DECL_P (t)
 	  && DECL_LOCAL_DECL_P (t)
@@ -48237,7 +48307,7 @@ cp_parser_omp_declare_target (cp_parser *parser, cp_token *pragma_tok)
 	  && DECL_LOCAL_DECL_ALIAS (t)
 	  && DECL_LOCAL_DECL_ALIAS (t) != error_mark_node)
 	handle_omp_declare_target_clause (c, DECL_LOCAL_DECL_ALIAS (t),
-					  device_type);
+					  device_type, indirect);
     }
   if (device_type && only_device_type)
     error_at (OMP_CLAUSE_LOCATION (clauses),
@@ -48250,7 +48320,8 @@ cp_parser_omp_declare_target (cp_parser *parser, cp_token *pragma_tok)
    # pragma omp begin declare target clauses[optseq] new-line  */
 
 #define OMP_BEGIN_DECLARE_TARGET_CLAUSE_MASK			\
-	(OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE_TYPE)
+	( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE_TYPE)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_INDIRECT))
 
 static void
 cp_parser_omp_begin (cp_parser *parser, cp_token *pragma_tok)
@@ -48280,11 +48351,16 @@ cp_parser_omp_begin (cp_parser *parser, cp_token *pragma_tok)
 					 "#pragma omp begin declare target",
 					 pragma_tok);
 	  int device_type = 0;
+	  bool indirect = 0;
 	  for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
-	    if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEVICE_TYPE)
-	      device_type |= OMP_CLAUSE_DEVICE_TYPE_KIND (c);
+	    {
+	      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEVICE_TYPE)
+		device_type |= OMP_CLAUSE_DEVICE_TYPE_KIND (c);
+	      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_INDIRECT)
+		indirect |= !integer_zerop (OMP_CLAUSE_INDIRECT_EXPR (c));
+	    }
 	  cp_omp_declare_target_attr a
-	    = { in_omp_attribute_pragma, device_type };
+	    = { in_omp_attribute_pragma, device_type, indirect };
 	  vec_safe_push (scope_chain->omp_declare_target_attribute, a);
 	}
       else
diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index 80ef1364e33..5fd7ccb187e 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -8830,6 +8830,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	case OMP_CLAUSE_IF_PRESENT:
 	case OMP_CLAUSE_FINALIZE:
 	case OMP_CLAUSE_NOHOST:
+	case OMP_CLAUSE_INDIRECT:
 	  break;
 
 	case OMP_CLAUSE_MERGEABLE:
diff --git a/gcc/lto-cgraph.cc b/gcc/lto-cgraph.cc
index 32c0f5ac6db..db6a22a444e 100644
--- a/gcc/lto-cgraph.cc
+++ b/gcc/lto-cgraph.cc
@@ -68,6 +68,7 @@ enum LTO_symtab_tags
   LTO_symtab_edge,
   LTO_symtab_indirect_edge,
   LTO_symtab_variable,
+  LTO_symtab_indirect_function,
   LTO_symtab_last_tag
 };
 
@@ -1111,6 +1112,18 @@ output_offload_tables (void)
 			       (*offload_vars)[i]);
     }
 
+  for (unsigned i = 0; i < vec_safe_length (offload_ind_funcs); i++)
+    {
+      symtab_node *node = symtab_node::get ((*offload_ind_funcs)[i]);
+      if (!node)
+	continue;
+      node->force_output = true;
+      streamer_write_enum (ob->main_stream, LTO_symtab_tags,
+			   LTO_symtab_last_tag, LTO_symtab_indirect_function);
+      lto_output_fn_decl_ref (ob->decl_state, ob->main_stream,
+			      (*offload_ind_funcs)[i]);
+    }
+
   if (output_requires)
     {
       HOST_WIDE_INT val = ((HOST_WIDE_INT) omp_requires_mask
@@ -1134,6 +1147,7 @@ output_offload_tables (void)
     {
       vec_free (offload_funcs);
       vec_free (offload_vars);
+      vec_free (offload_ind_funcs);
     }
 }
 
@@ -1863,6 +1877,19 @@ input_offload_tables (bool do_force_output)
 		varpool_node::get (var_decl)->force_output = 1;
 	      tmp_decl = var_decl;
 	    }
+	  else if (tag == LTO_symtab_indirect_function)
+	    {
+	      tree fn_decl
+		= lto_input_fn_decl_ref (ib, file_data);
+	      vec_safe_push (offload_ind_funcs, fn_decl);
+
+	      /* Prevent IPA from removing fn_decl as unreachable, since there
+		 may be no refs from the parent function to child_fn in offload
+		 LTO mode.  */
+	      if (do_force_output)
+		cgraph_node::get (fn_decl)->mark_force_output ();
+	      tmp_decl = fn_decl;
+	    }
 	  else if (tag == LTO_symtab_edge)
 	    {
 	      static bool error_emitted = false;
diff --git a/gcc/lto-section-names.h b/gcc/lto-section-names.h
index aa1b2f2eeff..f7ed622772f 100644
--- a/gcc/lto-section-names.h
+++ b/gcc/lto-section-names.h
@@ -37,5 +37,6 @@ extern const char *section_name_prefix;
 
 #define OFFLOAD_VAR_TABLE_SECTION_NAME ".gnu.offload_vars"
 #define OFFLOAD_FUNC_TABLE_SECTION_NAME ".gnu.offload_funcs"
+#define OFFLOAD_IND_FUNC_TABLE_SECTION_NAME ".gnu.offload_ind_funcs"
 
 #endif /* GCC_LTO_SECTION_NAMES_H */
diff --git a/gcc/omp-builtins.def b/gcc/omp-builtins.def
index e0f03263db0..ed78d49d205 100644
--- a/gcc/omp-builtins.def
+++ b/gcc/omp-builtins.def
@@ -445,6 +445,9 @@ DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_UPDATE, "GOMP_target_update_ext",
 DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_ENTER_EXIT_DATA,
 		  "GOMP_target_enter_exit_data",
 		  BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_UINT_PTR, ATTR_NOTHROW_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_MAP_INDIRECT_PTR,
+		  "GOMP_target_map_indirect_ptr",
+		  BT_FN_PTR_PTR, ATTR_NOTHROW_LIST)
 DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TEAMS4, "GOMP_teams4",
 		  BT_FN_BOOL_UINT_UINT_UINT_BOOL, ATTR_NOTHROW_LIST)
 DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TEAMS_REG, "GOMP_teams_reg",
diff --git a/gcc/omp-offload.cc b/gcc/omp-offload.cc
index 0d3c8794d54..3f14595b5b4 100644
--- a/gcc/omp-offload.cc
+++ b/gcc/omp-offload.cc
@@ -86,7 +86,7 @@ struct oacc_loop
 };
 
 /* Holds offload tables with decls.  */
-vec<tree, va_gc> *offload_funcs, *offload_vars;
+vec<tree, va_gc> *offload_funcs, *offload_vars, *offload_ind_funcs;
 
 /* Return level at which oacc routine may spawn a partitioned loop, or
    -1 if it is not a routine (i.e. is an offload fn).  */
@@ -351,6 +351,9 @@ omp_discover_implicit_declare_target (void)
     if (DECL_SAVED_TREE (node->decl))
       {
 	struct cgraph_node *cgn;
+	if (lookup_attribute ("omp declare target indirect",
+			      DECL_ATTRIBUTES (node->decl)))
+	  vec_safe_push (offload_ind_funcs, node->decl);
         if (omp_declare_target_fn_p (node->decl))
 	  worklist.safe_push (node->decl);
 	else if (DECL_STRUCT_FUNCTION (node->decl)
@@ -397,49 +400,66 @@ omp_finish_file (void)
 {
   unsigned num_funcs = vec_safe_length (offload_funcs);
   unsigned num_vars = vec_safe_length (offload_vars);
+  unsigned num_ind_funcs = vec_safe_length (offload_ind_funcs);
 
-  if (num_funcs == 0 && num_vars == 0)
+  if (num_funcs == 0 && num_vars == 0 && num_ind_funcs == 0)
     return;
 
   if (targetm_common.have_named_sections)
     {
-      vec<constructor_elt, va_gc> *v_f, *v_v;
+      vec<constructor_elt, va_gc> *v_f, *v_v, *v_if;
       vec_alloc (v_f, num_funcs);
       vec_alloc (v_v, num_vars * 2);
+      vec_alloc (v_if, num_ind_funcs);
 
       add_decls_addresses_to_decl_constructor (offload_funcs, v_f);
       add_decls_addresses_to_decl_constructor (offload_vars, v_v);
+      add_decls_addresses_to_decl_constructor (offload_ind_funcs, v_if);
 
       tree vars_decl_type = build_array_type_nelts (pointer_sized_int_node,
 						    vec_safe_length (v_v));
       tree funcs_decl_type = build_array_type_nelts (pointer_sized_int_node,
 						     num_funcs);
+      tree ind_funcs_decl_type = build_array_type_nelts (pointer_sized_int_node,
+							 num_ind_funcs);
+
       SET_TYPE_ALIGN (vars_decl_type, TYPE_ALIGN (pointer_sized_int_node));
       SET_TYPE_ALIGN (funcs_decl_type, TYPE_ALIGN (pointer_sized_int_node));
+      SET_TYPE_ALIGN (ind_funcs_decl_type, TYPE_ALIGN (pointer_sized_int_node));
       tree ctor_v = build_constructor (vars_decl_type, v_v);
       tree ctor_f = build_constructor (funcs_decl_type, v_f);
-      TREE_CONSTANT (ctor_v) = TREE_CONSTANT (ctor_f) = 1;
-      TREE_STATIC (ctor_v) = TREE_STATIC (ctor_f) = 1;
+      tree ctor_if = build_constructor (ind_funcs_decl_type, v_if);
+      TREE_CONSTANT (ctor_v) = TREE_CONSTANT (ctor_f) = TREE_CONSTANT (ctor_if) = 1;
+      TREE_STATIC (ctor_v) = TREE_STATIC (ctor_f) = TREE_STATIC (ctor_if) = 1;
       tree funcs_decl = build_decl (UNKNOWN_LOCATION, VAR_DECL,
 				    get_identifier (".offload_func_table"),
 				    funcs_decl_type);
       tree vars_decl = build_decl (UNKNOWN_LOCATION, VAR_DECL,
 				   get_identifier (".offload_var_table"),
 				   vars_decl_type);
-      TREE_STATIC (funcs_decl) = TREE_STATIC (vars_decl) = 1;
+      tree ind_funcs_decl = build_decl (UNKNOWN_LOCATION, VAR_DECL,
+					get_identifier (".offload_ind_func_table"),
+					ind_funcs_decl_type);
+      TREE_STATIC (funcs_decl) = TREE_STATIC (ind_funcs_decl) = 1;
+      TREE_STATIC (vars_decl) = 1;
       /* Do not align tables more than TYPE_ALIGN (pointer_sized_int_node),
 	 otherwise a joint table in a binary will contain padding between
 	 tables from multiple object files.  */
-      DECL_USER_ALIGN (funcs_decl) = DECL_USER_ALIGN (vars_decl) = 1;
+      DECL_USER_ALIGN (funcs_decl) = DECL_USER_ALIGN (ind_funcs_decl) = 1;
+      DECL_USER_ALIGN (vars_decl) = 1;
       SET_DECL_ALIGN (funcs_decl, TYPE_ALIGN (funcs_decl_type));
       SET_DECL_ALIGN (vars_decl, TYPE_ALIGN (vars_decl_type));
+      SET_DECL_ALIGN (ind_funcs_decl, TYPE_ALIGN (ind_funcs_decl_type));
       DECL_INITIAL (funcs_decl) = ctor_f;
       DECL_INITIAL (vars_decl) = ctor_v;
+      DECL_INITIAL (ind_funcs_decl) = ctor_if;
       set_decl_section_name (funcs_decl, OFFLOAD_FUNC_TABLE_SECTION_NAME);
       set_decl_section_name (vars_decl, OFFLOAD_VAR_TABLE_SECTION_NAME);
-
+      set_decl_section_name (ind_funcs_decl,
+			     OFFLOAD_IND_FUNC_TABLE_SECTION_NAME);
       varpool_node::finalize_decl (vars_decl);
       varpool_node::finalize_decl (funcs_decl);
+      varpool_node::finalize_decl (ind_funcs_decl);
     }
   else
     {
@@ -471,6 +491,15 @@ omp_finish_file (void)
 #endif
 	    targetm.record_offload_symbol (it);
 	}
+      for (unsigned i = 0; i < num_ind_funcs; i++)
+	{
+	  tree it = (*offload_ind_funcs)[i];
+	  /* See also add_decls_addresses_to_decl_constructor
+	     and output_offload_tables in lto-cgraph.cc.  */
+	  if (!in_lto_p && !symtab_node::get (it))
+	    continue;
+	  targetm.record_offload_symbol (it);
+	}
     }
 }
 
@@ -2603,6 +2632,8 @@ execute_omp_device_lower ()
   gimple_stmt_iterator gsi;
   bool calls_declare_variant_alt
     = cgraph_node::get (cfun->decl)->calls_declare_variant_alt;
+  bool omp_redirect_indirect_calls
+    = vec_safe_length (offload_ind_funcs) > 0;
   FOR_EACH_BB_FN (bb, cfun)
     for (gsi = gsi_start_bb (bb); !gsi_end_p (gsi); gsi_next (&gsi))
       {
@@ -2621,6 +2652,33 @@ execute_omp_device_lower ()
 		      update_stmt (stmt);
 		    }
 		}
+	    if (omp_redirect_indirect_calls
+		&& gimple_call_fndecl (stmt) == NULL_TREE)
+	      {
+		gcall *orig_call = dyn_cast <gcall *> (stmt);
+		tree call_fn = gimple_call_fn (stmt);
+		tree map_ptr_fn
+		  = builtin_decl_explicit (BUILT_IN_GOMP_TARGET_MAP_INDIRECT_PTR);
+		tree fn_ty = TREE_TYPE (call_fn);
+
+		if (TREE_CODE (call_fn) == OBJ_TYPE_REF)
+		  {
+		    tree obj_ref = create_tmp_reg (TREE_TYPE (call_fn),
+						   ".ind_fn_objref");
+		    gimple *gassign = gimple_build_assign (obj_ref, call_fn);
+		    gsi_insert_before  (&gsi, gassign, GSI_SAME_STMT);
+		    call_fn = obj_ref;
+		  }
+		tree mapped_fn = create_tmp_reg (fn_ty, ".ind_fn");
+		gimple *gcall =
+		    gimple_build_call (map_ptr_fn, 1, call_fn);
+		gimple_set_location (gcall, gimple_location (stmt));
+		gimple_call_set_lhs (gcall, mapped_fn);
+		gsi_insert_before (&gsi, gcall, GSI_SAME_STMT);
+
+		gimple_call_set_fn (orig_call, mapped_fn);
+		update_stmt (orig_call);
+	      }
 	    continue;
 	  }
 	tree lhs = gimple_call_lhs (stmt), rhs = NULL_TREE;
@@ -2759,7 +2817,8 @@ public:
     {
       return (!(fun->curr_properties & PROP_gimple_lomp_dev)
 	      || (flag_openmp
-		  && cgraph_node::get (fun->decl)->calls_declare_variant_alt));
+		  && (cgraph_node::get (fun->decl)->calls_declare_variant_alt
+		      || vec_safe_length (offload_ind_funcs) > 0)));
     }
   unsigned int execute (function *) final override
     {
diff --git a/gcc/omp-offload.h b/gcc/omp-offload.h
index 73711e74c7d..ae364422417 100644
--- a/gcc/omp-offload.h
+++ b/gcc/omp-offload.h
@@ -28,6 +28,7 @@ extern int oacc_fn_attrib_level (tree attr);
 
 extern GTY(()) vec<tree, va_gc> *offload_funcs;
 extern GTY(()) vec<tree, va_gc> *offload_vars;
+extern GTY(()) vec<tree, va_gc> *offload_ind_funcs;
 
 extern void omp_finish_file (void);
 extern void omp_discover_implicit_declare_target (void);
diff --git a/gcc/testsuite/c-c++-common/gomp/declare-target-indirect-1.c b/gcc/testsuite/c-c++-common/gomp/declare-target-indirect-1.c
new file mode 100644
index 00000000000..45e3ee598da
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/declare-target-indirect-1.c
@@ -0,0 +1,51 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp" } */
+
+extern int a, b;
+#define X 1
+#define Y 0
+
+#pragma omp begin declare target indirect
+void fn1 (void) { }
+#pragma omp end declare target
+
+#pragma omp begin declare target indirect (1)
+void fn2 (void) { }
+#pragma omp end declare target
+
+#pragma omp begin declare target indirect (0)
+void fn3 (void) { }
+#pragma omp end declare target
+
+void fn4 (void) { }
+#pragma omp declare target indirect to (fn4)
+
+void fn5 (void) { }
+#pragma omp declare target indirect (1) to (fn5)
+
+void fn6 (void) { }
+#pragma omp declare target indirect (0) to (fn6)
+
+void fn7 (void) { }
+#pragma omp declare target indirect (-1) to (fn7)
+
+/* Compile-time non-constant expressions are not allowed.  */
+void fn8 (void) { }
+#pragma omp declare target indirect (a + b) to (fn8) /* { dg-error "expected constant integer expression" } */
+
+/* Compile-time constant expressions are permissible.  */
+void fn9 (void) { }
+#pragma omp declare target indirect (X*Y) to (fn9)
+
+/* 'omp declare target'...'omp end declare target' form cannot take
+   clauses.  */
+#pragma omp declare target indirect
+void fn10 (void) { }
+#pragma omp end declare target /* { dg-error "'#pragma omp end declare target' without corresponding '#pragma omp declare target' or '#pragma omp begin declare target'" } */
+
+void fn11 (void) { }
+#pragma omp declare target indirect (1) indirect (0) to (fn11) /* { dg-error "too many .indirect. clauses" } */
+
+/* Indirect on a variable should have no effect.  */
+int x;
+#pragma omp declare target indirect to(x)
diff --git a/gcc/testsuite/c-c++-common/gomp/declare-target-indirect-2.c b/gcc/testsuite/c-c++-common/gomp/declare-target-indirect-2.c
new file mode 100644
index 00000000000..48ba4f86362
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/declare-target-indirect-2.c
@@ -0,0 +1,20 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp -fdump-tree-gimple" } */
+
+#pragma omp begin declare target indirect
+void fn1 (void) { }
+#pragma omp end declare target
+/* { dg-final { scan-tree-dump "__attribute__\\\(\\\(omp declare target, omp declare target block, omp declare target indirect\\\)\\\)\\\nvoid fn1" "gimple" } } */
+
+#pragma omp begin declare target indirect (0)
+void fn2 (void) { }
+#pragma omp end declare target
+/* { dg-final { scan-tree-dump "__attribute__\\\(\\\(omp declare target, omp declare target block\\\)\\\)\\\nvoid fn2" "gimple" } } */
+
+void fn3 (void) { }
+#pragma omp declare target indirect to (fn3)
+/* { dg-final { scan-tree-dump "__attribute__\\\(\\\(omp declare target indirect, omp declare target\\\)\\\)\\\nvoid fn3" "gimple" } } */
+
+void fn4 (void) { }
+#pragma omp declare target indirect (0) to (fn4)
+/* { dg-final { scan-tree-dump "__attribute__\\\(\\\(omp declare target\\\)\\\)\\\nvoid fn4" "gimple" } } */
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index 91551fde900..f7549017df6 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -347,6 +347,9 @@ enum omp_clause_code {
   /* OpenMP clause: doacross ({source,sink}:vec).  */
   OMP_CLAUSE_DOACROSS,
 
+  /* OpenMP clause: indirect [(constant-integer-expression)].  */
+  OMP_CLAUSE_INDIRECT,
+
   /* Internal structure to hold OpenACC cache directive's variable-list.
      #pragma acc cache (variable-list).  */
   OMP_CLAUSE__CACHE_,
diff --git a/gcc/tree.cc b/gcc/tree.cc
index 54ca5e750df..34a03ff04c8 100644
--- a/gcc/tree.cc
+++ b/gcc/tree.cc
@@ -269,6 +269,7 @@ unsigned const char omp_clause_num_ops[] =
   2, /* OMP_CLAUSE_MAP  */
   1, /* OMP_CLAUSE_HAS_DEVICE_ADDR  */
   1, /* OMP_CLAUSE_DOACROSS  */
+  1, /* OMP_CLAUSE_INDIRECT  */
   2, /* OMP_CLAUSE__CACHE_  */
   2, /* OMP_CLAUSE_GANG  */
   1, /* OMP_CLAUSE_ASYNC  */
@@ -360,6 +361,7 @@ const char * const omp_clause_code_name[] =
   "map",
   "has_device_addr",
   "doacross",
+  "indirect",
   "_cache_",
   "gang",
   "async",
diff --git a/gcc/tree.h b/gcc/tree.h
index 005c157e9b0..78b15c24641 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1842,6 +1842,10 @@ class auto_suppress_location_wrappers
 #define OMP_CLAUSE_DEVICE_TYPE_KIND(NODE) \
   (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_DEVICE_TYPE)->omp_clause.subcode.device_type_kind)
 
+#define OMP_CLAUSE_INDIRECT_EXPR(NODE) \
+  OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_INDIRECT), 0)
+
+
 /* True if there is a device clause with a device-modifier 'ancestor'.  */
 #define OMP_CLAUSE_DEVICE_ANCESTOR(NODE) \
   (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_DEVICE)->base.public_flag)
diff --git a/libgcc/offloadstuff.c b/libgcc/offloadstuff.c
index 4e1c4d41dd5..18c5bf89b69 100644
--- a/libgcc/offloadstuff.c
+++ b/libgcc/offloadstuff.c
@@ -43,6 +43,7 @@ see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
 #if defined(HAVE_GAS_HIDDEN) && ENABLE_OFFLOADING == 1
 
 #define OFFLOAD_FUNC_TABLE_SECTION_NAME ".gnu.offload_funcs"
+#define OFFLOAD_IND_FUNC_TABLE_SECTION_NAME ".gnu.offload_ind_funcs"
 #define OFFLOAD_VAR_TABLE_SECTION_NAME ".gnu.offload_vars"
 
 #ifdef CRT_BEGIN
@@ -53,6 +54,9 @@ const void *const __offload_func_table[0]
 const void *const __offload_var_table[0]
   __attribute__ ((__used__, visibility ("hidden"),
 		  section (OFFLOAD_VAR_TABLE_SECTION_NAME))) = { };
+const void *const __offload_ind_func_table[0]
+  __attribute__ ((__used__, visibility ("hidden"),
+		  section (OFFLOAD_IND_FUNC_TABLE_SECTION_NAME))) = { };
 
 #elif defined CRT_END
 
@@ -62,19 +66,25 @@ const void *const __offload_funcs_end[0]
 const void *const __offload_vars_end[0]
   __attribute__ ((__used__, visibility ("hidden"),
 		  section (OFFLOAD_VAR_TABLE_SECTION_NAME))) = { };
+const void *const __offload_ind_funcs_end[0]
+  __attribute__ ((__used__, visibility ("hidden"),
+		  section (OFFLOAD_IND_FUNC_TABLE_SECTION_NAME))) = { };
 
 #elif defined CRT_TABLE
 
 extern const void *const __offload_func_table[];
 extern const void *const __offload_var_table[];
+extern const void *const __offload_ind_func_table[];
 extern const void *const __offload_funcs_end[];
 extern const void *const __offload_vars_end[];
+extern const void *const __offload_ind_funcs_end[];
 
 const void *const __OFFLOAD_TABLE__[]
   __attribute__ ((__visibility__ ("hidden"))) =
 {
   &__offload_func_table, &__offload_funcs_end,
-  &__offload_var_table, &__offload_vars_end
+  &__offload_var_table, &__offload_vars_end,
+  &__offload_ind_func_table, &__offload_ind_funcs_end,
 };
 
 #else /* ! CRT_BEGIN && ! CRT_END && ! CRT_TABLE  */
diff --git a/libgomp/Makefile.am b/libgomp/Makefile.am
index 428f7a9dab5..2549ea8906b 100644
--- a/libgomp/Makefile.am
+++ b/libgomp/Makefile.am
@@ -67,7 +67,7 @@ libgomp_la_SOURCES = alloc.c atomic.c barrier.c critical.c env.c error.c \
 	target.c splay-tree.c libgomp-plugin.c oacc-parallel.c oacc-host.c \
 	oacc-init.c oacc-mem.c oacc-async.c oacc-plugin.c oacc-cuda.c \
 	priority_queue.c affinity-fmt.c teams.c allocator.c oacc-profiling.c \
-	oacc-target.c
+	oacc-target.c target-indirect.c
 
 include $(top_srcdir)/plugin/Makefrag.am
 
diff --git a/libgomp/Makefile.in b/libgomp/Makefile.in
index 431bc87b629..32e64292a6d 100644
--- a/libgomp/Makefile.in
+++ b/libgomp/Makefile.in
@@ -219,7 +219,7 @@ am_libgomp_la_OBJECTS = alloc.lo atomic.lo barrier.lo critical.lo \
 	oacc-parallel.lo oacc-host.lo oacc-init.lo oacc-mem.lo \
 	oacc-async.lo oacc-plugin.lo oacc-cuda.lo priority_queue.lo \
 	affinity-fmt.lo teams.lo allocator.lo oacc-profiling.lo \
-	oacc-target.lo $(am__objects_1)
+	oacc-target.lo target-indirect.lo $(am__objects_1)
 libgomp_la_OBJECTS = $(am_libgomp_la_OBJECTS)
 AM_V_P = $(am__v_P_@AM_V@)
 am__v_P_ = $(am__v_P_@AM_DEFAULT_V@)
@@ -549,7 +549,7 @@ libgomp_la_SOURCES = alloc.c atomic.c barrier.c critical.c env.c \
 	oacc-parallel.c oacc-host.c oacc-init.c oacc-mem.c \
 	oacc-async.c oacc-plugin.c oacc-cuda.c priority_queue.c \
 	affinity-fmt.c teams.c allocator.c oacc-profiling.c \
-	oacc-target.c $(am__append_3)
+	oacc-target.c target-indirect.c $(am__append_3)
 
 # Nvidia PTX OpenACC plugin.
 @PLUGIN_NVPTX_TRUE@libgomp_plugin_nvptx_version_info = -version-info $(libtool_VERSION)
@@ -777,6 +777,7 @@ distclean-compile:
 @AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/sem.Plo@am__quote@
 @AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/single.Plo@am__quote@
 @AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/splay-tree.Plo@am__quote@
+@AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/target-indirect.Plo@am__quote@
 @AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/target.Plo@am__quote@
 @AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/task.Plo@am__quote@
 @AMDEP_TRUE@@am__include@ @am__quote@./$(DEPDIR)/team.Plo@am__quote@
diff --git a/libgomp/config/accel/target-indirect.c b/libgomp/config/accel/target-indirect.c
new file mode 100644
index 00000000000..6ee82a0ebd0
--- /dev/null
+++ b/libgomp/config/accel/target-indirect.c
@@ -0,0 +1,126 @@
+/* Copyright (C) 2023 Free Software Foundation, Inc.
+
+   Contributed by Siemens.
+
+   This file is part of the GNU Offloading and Multi Processing Library
+   (libgomp).
+
+   Libgomp is free software; you can redistribute it and/or modify it
+   under the terms of the GNU General Public License as published by
+   the Free Software Foundation; either version 3, or (at your option)
+   any later version.
+
+   Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+   WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+   FOR A PARTICULAR PURPOSE.  See the GNU General Public License for
+   more details.
+
+   Under Section 7 of GPL version 3, you are granted additional
+   permissions described in the GCC Runtime Library Exception, version
+   3.1, as published by the Free Software Foundation.
+
+   You should have received a copy of the GNU General Public License and
+   a copy of the GCC Runtime Library Exception along with this program;
+   see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
+   <http://www.gnu.org/licenses/>.  */
+
+#include <assert.h>
+#include "libgomp.h"
+
+#define splay_tree_prefix indirect
+#define splay_tree_c
+#include "splay-tree.h"
+
+volatile void **GOMP_INDIRECT_ADDR_MAP = NULL;
+
+/* Use a splay tree to lookup the target address instead of using a
+   linear search.  */
+#define USE_SPLAY_TREE_LOOKUP
+
+#ifdef USE_SPLAY_TREE_LOOKUP
+
+static struct indirect_splay_tree_s indirect_map;
+static indirect_splay_tree_node indirect_array = NULL;
+
+/* Build the splay tree used for host->target address lookups.  */
+
+void
+build_indirect_map (void)
+{
+  size_t num_ind_funcs = 0;
+  volatile void **map_entry;
+  static int lock = 0; /* == gomp_mutex_t lock; gomp_mutex_init (&lock); */
+
+  if (!GOMP_INDIRECT_ADDR_MAP)
+    return;
+
+  gomp_mutex_lock (&lock);
+
+  if (!indirect_array)
+    {
+      /* Count the number of entries in the NULL-terminated address map.  */
+      for (map_entry = GOMP_INDIRECT_ADDR_MAP; *map_entry;
+	   map_entry += 2, num_ind_funcs++);
+
+      /* Build splay tree for address lookup.  */
+      indirect_array = gomp_malloc (num_ind_funcs * sizeof (*indirect_array));
+      indirect_splay_tree_node array = indirect_array;
+      map_entry = GOMP_INDIRECT_ADDR_MAP;
+
+      for (int i = 0; i < num_ind_funcs; i++, array++)
+	{
+	  indirect_splay_tree_key k = &array->key;
+	  k->host_addr = (uint64_t) *map_entry++;
+	  k->target_addr = (uint64_t) *map_entry++;
+	  array->left = NULL;
+	  array->right = NULL;
+	  indirect_splay_tree_insert (&indirect_map, array);
+	}
+    }
+
+  gomp_mutex_unlock (&lock);
+}
+
+void *
+GOMP_target_map_indirect_ptr (void *ptr)
+{
+  /* NULL pointers always resolve to NULL.  */
+  if (!ptr)
+    return ptr;
+
+  assert (indirect_array);
+
+  struct indirect_splay_tree_key_s k;
+  indirect_splay_tree_key node = NULL;
+
+  k.host_addr = (uint64_t) ptr;
+  node = indirect_splay_tree_lookup (&indirect_map, &k);
+
+  return node ? (void *) node->target_addr : ptr;
+}
+
+#else
+
+void
+build_indirect_map (void)
+{
+}
+
+void *
+GOMP_target_map_indirect_ptr (void *ptr)
+{
+  /* NULL pointers always resolve to NULL.  */
+  if (!ptr)
+    return ptr;
+
+  assert (GOMP_INDIRECT_ADDR_MAP);
+
+  for (volatile void **map_entry = GOMP_INDIRECT_ADDR_MAP; *map_entry;
+       map_entry += 2)
+    if (*map_entry == ptr)
+      return (void *) *(map_entry + 1);
+
+  return ptr;
+}
+
+#endif
diff --git a/libgomp/config/gcn/team.c b/libgomp/config/gcn/team.c
index f03207c84e3..fb20cbbcf9f 100644
--- a/libgomp/config/gcn/team.c
+++ b/libgomp/config/gcn/team.c
@@ -30,6 +30,7 @@
 #include <string.h>
 
 static void gomp_thread_start (struct gomp_thread_pool *);
+extern void build_indirect_map (void);
 
 /* This externally visible function handles target region entry.  It
    sets up a per-team thread pool and transfers control by returning to
@@ -45,6 +46,9 @@ gomp_gcn_enter_kernel (void)
 {
   int threadid = __builtin_gcn_dim_pos (1);
 
+  /* Initialize indirect function support.  */
+  build_indirect_map ();
+
   if (threadid == 0)
     {
       int numthreads = __builtin_gcn_dim_size (1);
diff --git a/libgomp/config/linux/target-indirect.c b/libgomp/config/linux/target-indirect.c
new file mode 100644
index 00000000000..de2c708bb81
--- /dev/null
+++ b/libgomp/config/linux/target-indirect.c
@@ -0,0 +1,31 @@
+/* Copyright (C) 2023 Free Software Foundation, Inc.
+
+   Contributed by Siemens.
+
+   This file is part of the GNU Offloading and Multi Processing Library
+   (libgomp).
+
+   Libgomp is free software; you can redistribute it and/or modify it
+   under the terms of the GNU General Public License as published by
+   the Free Software Foundation; either version 3, or (at your option)
+   any later version.
+
+   Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+   WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+   FOR A PARTICULAR PURPOSE.  See the GNU General Public License for
+   more details.
+
+   Under Section 7 of GPL version 3, you are granted additional
+   permissions described in the GCC Runtime Library Exception, version
+   3.1, as published by the Free Software Foundation.
+
+   You should have received a copy of the GNU General Public License and
+   a copy of the GCC Runtime Library Exception along with this program;
+   see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
+   <http://www.gnu.org/licenses/>.  */
+
+void *
+GOMP_target_map_indirect_ptr (void *ptr)
+{
+  return ptr;
+}
diff --git a/libgomp/config/nvptx/team.c b/libgomp/config/nvptx/team.c
index af5f3171a47..59521fabd99 100644
--- a/libgomp/config/nvptx/team.c
+++ b/libgomp/config/nvptx/team.c
@@ -35,6 +35,7 @@ struct gomp_thread *nvptx_thrs __attribute__((shared,nocommon));
 int __gomp_team_num __attribute__((shared,nocommon));
 
 static void gomp_thread_start (struct gomp_thread_pool *);
+extern void build_indirect_map (void);
 
 
 /* This externally visible function handles target region entry.  It
@@ -52,6 +53,10 @@ gomp_nvptx_main (void (*fn) (void *), void *fn_data)
   int tid, ntids;
   asm ("mov.u32 %0, %%tid.y;" : "=r" (tid));
   asm ("mov.u32 %0, %%ntid.y;" : "=r" (ntids));
+
+  /* Initialize indirect function support.  */
+  build_indirect_map ();
+
   if (tid == 0)
     {
       gomp_global_icv.nthreads_var = ntids;
diff --git a/libgomp/libgomp-plugin.h b/libgomp/libgomp-plugin.h
index dc993882c3b..3ce032c5cc0 100644
--- a/libgomp/libgomp-plugin.h
+++ b/libgomp/libgomp-plugin.h
@@ -107,6 +107,8 @@ struct addr_pair
    must be stringified).  */
 #define GOMP_ADDITIONAL_ICVS __gomp_additional_icvs
 
+#define GOMP_INDIRECT_ADDR_MAP __gomp_indirect_addr_map
+
 /* Miscellaneous functions.  */
 extern void *GOMP_PLUGIN_malloc (size_t) __attribute__ ((malloc));
 extern void *GOMP_PLUGIN_malloc_cleared (size_t) __attribute__ ((malloc));
@@ -132,7 +134,8 @@ extern bool GOMP_OFFLOAD_init_device (int);
 extern bool GOMP_OFFLOAD_fini_device (int);
 extern unsigned GOMP_OFFLOAD_version (void);
 extern int GOMP_OFFLOAD_load_image (int, unsigned, const void *,
-				    struct addr_pair **, uint64_t **);
+				    struct addr_pair **, uint64_t **,
+				    uint64_t *);
 extern bool GOMP_OFFLOAD_unload_image (int, unsigned, const void *);
 extern void *GOMP_OFFLOAD_alloc (int, size_t);
 extern bool GOMP_OFFLOAD_free (int, void *);
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index 68f20651fbf..15a767cf317 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -1274,6 +1274,29 @@ reverse_splay_compare (reverse_splay_tree_key x, reverse_splay_tree_key y)
 #define splay_tree_prefix reverse
 #include "splay-tree.h"
 
+/* Indirect target function splay-tree handling.  */
+
+struct indirect_splay_tree_key_s {
+  uint64_t host_addr, target_addr;
+};
+
+typedef struct indirect_splay_tree_node_s *indirect_splay_tree_node;
+typedef struct indirect_splay_tree_s *indirect_splay_tree;
+typedef struct indirect_splay_tree_key_s *indirect_splay_tree_key;
+
+static inline int
+indirect_splay_compare (indirect_splay_tree_key x, indirect_splay_tree_key y)
+{
+  if (x->host_addr < y->host_addr)
+    return -1;
+  if (x->host_addr > y->host_addr)
+    return 1;
+  return 0;
+}
+
+#define splay_tree_prefix indirect
+#include "splay-tree.h"
+
 struct target_mem_desc {
   /* Reference count.  */
   uintptr_t refcount;
diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map
index ce6b719a57f..90c401453b2 100644
--- a/libgomp/libgomp.map
+++ b/libgomp/libgomp.map
@@ -419,6 +419,7 @@ GOMP_5.1 {
 GOMP_5.1.1 {
   global:
 	GOMP_taskwait_depend_nowait;
+	GOMP_target_map_indirect_ptr;
 } GOMP_5.1;
 
 OACC_2.0 {
diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi
index d24f590fd84..cd2ef065456 100644
--- a/libgomp/libgomp.texi
+++ b/libgomp/libgomp.texi
@@ -304,7 +304,7 @@ The OpenMP 4.5 specification is fully supported.
 @item Iterators in @code{target update} motion clauses and @code{map}
       clauses @tab N @tab
 @item Indirect calls to the device version of a procedure or function in
-      @code{target} regions @tab N @tab
+      @code{target} regions @tab P @tab Only C and C++
 @item @code{interop} directive @tab N @tab
 @item @code{omp_interop_t} object support in runtime routines @tab N @tab
 @item @code{nowait} clause in @code{taskwait} directive @tab Y @tab
@@ -353,7 +353,7 @@ to address of matching mapped list item per 5.1, Sect. 2.21.7.2 @tab N @tab
 @item For Fortran, diagnose placing declarative before/between @code{USE},
       @code{IMPORT}, and @code{IMPLICIT} as invalid @tab N @tab
 @item Optional comma between directive and clause in the @code{#pragma} form @tab Y @tab
-@item @code{indirect} clause in @code{declare target} @tab N @tab
+@item @code{indirect} clause in @code{declare target} @tab P @tab Only C and C++
 @item @code{device_type(nohost)}/@code{device_type(host)} for variables @tab N @tab
 @item @code{present} modifier to the @code{map}, @code{to} and @code{from}
       clauses @tab Y @tab
diff --git a/libgomp/libgomp_g.h b/libgomp/libgomp_g.h
index 5c1675c7869..95046312ae9 100644
--- a/libgomp/libgomp_g.h
+++ b/libgomp/libgomp_g.h
@@ -357,6 +357,7 @@ extern void GOMP_target_enter_exit_data (int, size_t, void **, size_t *,
 					 void **);
 extern void GOMP_teams (unsigned int, unsigned int);
 extern bool GOMP_teams4 (unsigned int, unsigned int, unsigned int, bool);
+extern void *GOMP_target_map_indirect_ptr (void *);
 
 /* teams.c */
 
diff --git a/libgomp/oacc-host.c b/libgomp/oacc-host.c
index 5980d510838..fbab75d7d43 100644
--- a/libgomp/oacc-host.c
+++ b/libgomp/oacc-host.c
@@ -82,7 +82,8 @@ host_load_image (int n __attribute__ ((unused)),
 		 unsigned v __attribute__ ((unused)),
 		 const void *t __attribute__ ((unused)),
 		 struct addr_pair **r __attribute__ ((unused)),
-		 uint64_t **f __attribute__ ((unused)))
+		 uint64_t **f __attribute__ ((unused)),
+		 uint64_t *i __attribute__ ((unused)))
 {
   return 0;
 }
diff --git a/libgomp/plugin/plugin-gcn.c b/libgomp/plugin/plugin-gcn.c
index ef22d48da79..597bd75798d 100644
--- a/libgomp/plugin/plugin-gcn.c
+++ b/libgomp/plugin/plugin-gcn.c
@@ -365,6 +365,7 @@ struct gcn_image_desc
   } *gcn_image;
   const unsigned kernel_count;
   struct hsa_kernel_description *kernel_infos;
+  const unsigned ind_func_count;
   const unsigned global_variable_count;
 };
 
@@ -3359,7 +3360,8 @@ GOMP_OFFLOAD_init_device (int n)
 int
 GOMP_OFFLOAD_load_image (int ord, unsigned version, const void *target_data,
 			 struct addr_pair **target_table,
-			 uint64_t **rev_fn_table)
+			 uint64_t **rev_fn_table,
+			 uint64_t *host_ind_fn_table)
 {
   if (GOMP_VERSION_DEV (version) != GOMP_VERSION_GCN)
     {
@@ -3375,6 +3377,7 @@ GOMP_OFFLOAD_load_image (int ord, unsigned version, const void *target_data,
   struct module_info *module;
   struct kernel_info *kernel;
   int kernel_count = image_desc->kernel_count;
+  unsigned ind_func_count = image_desc->ind_func_count;
   unsigned var_count = image_desc->global_variable_count;
   /* Currently, "others" is a struct of ICVS.  */
   int other_count = 1;
@@ -3393,6 +3396,7 @@ GOMP_OFFLOAD_load_image (int ord, unsigned version, const void *target_data,
     return -1;
 
   GCN_DEBUG ("Encountered %d kernels in an image\n", kernel_count);
+  GCN_DEBUG ("Encountered %d indirect functions in an image\n", ind_func_count);
   GCN_DEBUG ("Encountered %u global variables in an image\n", var_count);
   GCN_DEBUG ("Expect %d other variables in an image\n", other_count);
   pair = GOMP_PLUGIN_malloc ((kernel_count + var_count + other_count - 2)
@@ -3474,6 +3478,87 @@ GOMP_OFFLOAD_load_image (int ord, unsigned version, const void *target_data,
 	}
     }
 
+  if (ind_func_count > 0)
+    {
+      hsa_status_t status;
+
+      /* Read indirect function table from image.  */
+      hsa_executable_symbol_t ind_funcs_symbol;
+      status = hsa_fns.hsa_executable_get_symbol_fn (agent->executable, NULL,
+						     ".offload_ind_func_table",
+						     agent->id,
+						     0, &ind_funcs_symbol);
+
+      if (status != HSA_STATUS_SUCCESS)
+	hsa_fatal ("Could not find .offload_ind_func_table symbol in the "
+		   "code object", status);
+
+      uint64_t ind_funcs_table_addr;
+      status = hsa_fns.hsa_executable_symbol_get_info_fn
+	(ind_funcs_symbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS,
+	 &ind_funcs_table_addr);
+      if (status != HSA_STATUS_SUCCESS)
+	hsa_fatal ("Could not extract a variable from its symbol", status);
+
+      uint64_t ind_funcs_table[ind_func_count];
+      GOMP_OFFLOAD_dev2host (agent->device_id, ind_funcs_table,
+			     (void*) ind_funcs_table_addr,
+			     sizeof (ind_funcs_table));
+
+      /* Build host->target address map for indirect functions.  */
+      uint64_t ind_fn_map[ind_func_count * 2 + 1];
+      for (unsigned i = 0; i < ind_func_count; i++)
+	{
+	  ind_fn_map[i * 2] = host_ind_fn_table[i];
+	  ind_fn_map[i * 2 + 1] = ind_funcs_table[i];
+	  GCN_DEBUG ("Indirect function %d: %lx->%lx\n",
+		     i, host_ind_fn_table[i], ind_funcs_table[i]);
+	}
+      ind_fn_map[ind_func_count * 2] = 0;
+
+      /* Write the map onto the target.  */
+      void *map_target_addr
+	= GOMP_OFFLOAD_alloc (agent->device_id, sizeof (ind_fn_map));
+      GCN_DEBUG ("Allocated indirect map at %p\n", map_target_addr);
+
+      GOMP_OFFLOAD_host2dev (agent->device_id, map_target_addr,
+			     (void*) ind_fn_map,
+			     sizeof (ind_fn_map));
+
+      /* Write address of the map onto the target.  */
+      hsa_executable_symbol_t symbol;
+
+      status
+	= hsa_fns.hsa_executable_get_symbol_fn (agent->executable, NULL,
+						XSTRING (GOMP_INDIRECT_ADDR_MAP),
+						agent->id, 0, &symbol);
+      if (status != HSA_STATUS_SUCCESS)
+	hsa_fatal ("Could not find GOMP_INDIRECT_ADDR_MAP in code object",
+		   status);
+
+      uint64_t varptr;
+      uint32_t varsize;
+
+      status = hsa_fns.hsa_executable_symbol_get_info_fn
+	(symbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS,
+	 &varptr);
+      if (status != HSA_STATUS_SUCCESS)
+	hsa_fatal ("Could not extract a variable from its symbol", status);
+      status = hsa_fns.hsa_executable_symbol_get_info_fn
+	(symbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SIZE,
+	&varsize);
+      if (status != HSA_STATUS_SUCCESS)
+	hsa_fatal ("Could not extract a variable size from its symbol",
+		   status);
+
+      GCN_DEBUG ("Found GOMP_INDIRECT_ADDR_MAP at %lx with size %d\n",
+		 varptr, varsize);
+
+      GOMP_OFFLOAD_host2dev (agent->device_id, (void *) varptr,
+			     &map_target_addr,
+			     sizeof (map_target_addr));
+    }
+
   GCN_DEBUG ("Looking for variable %s\n", XSTRING (GOMP_ADDITIONAL_ICVS));
 
   hsa_status_t status;
diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index 00d4241ae02..78aa95c843a 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -266,6 +266,8 @@ typedef struct nvptx_tdata
 
   const struct targ_fn_launch *fn_descs;
   unsigned fn_num;
+
+  unsigned ind_fn_num;
 } nvptx_tdata_t;
 
 /* Descriptor of a loaded function.  */
@@ -1285,12 +1287,13 @@ nvptx_set_clocktick (CUmodule module, struct ptx_device *dev)
 int
 GOMP_OFFLOAD_load_image (int ord, unsigned version, const void *target_data,
 			 struct addr_pair **target_table,
-			 uint64_t **rev_fn_table)
+			 uint64_t **rev_fn_table,
+			 uint64_t *host_ind_fn_table)
 {
   CUmodule module;
   const char *const *var_names;
   const struct targ_fn_launch *fn_descs;
-  unsigned int fn_entries, var_entries, other_entries, i, j;
+  unsigned int fn_entries, var_entries, ind_fn_entries, other_entries, i, j;
   struct targ_fn_descriptor *targ_fns;
   struct addr_pair *targ_tbl;
   const nvptx_tdata_t *img_header = (const nvptx_tdata_t *) target_data;
@@ -1319,6 +1322,7 @@ GOMP_OFFLOAD_load_image (int ord, unsigned version, const void *target_data,
   var_names = img_header->var_names;
   fn_entries = img_header->fn_num;
   fn_descs = img_header->fn_descs;
+  ind_fn_entries = img_header->ind_fn_num;
 
   /* Currently, other_entries contains only the struct of ICVs.  */
   other_entries = 1;
@@ -1373,6 +1377,60 @@ GOMP_OFFLOAD_load_image (int ord, unsigned version, const void *target_data,
       targ_tbl->end = targ_tbl->start + bytes;
     }
 
+  if (ind_fn_entries > 0)
+    {
+      CUdeviceptr var;
+      size_t bytes;
+
+      /* Read indirect function table from image.  */
+      CUresult r = CUDA_CALL_NOCHECK (cuModuleGetGlobal, &var, &bytes, module,
+				      "$offload_ind_func_table");
+      if (r != CUDA_SUCCESS)
+	GOMP_PLUGIN_fatal ("cuModuleGetGlobal error: %s", cuda_error (r));
+      assert (bytes == sizeof (uint64_t) * ind_fn_entries);
+
+      uint64_t ind_fn_table[ind_fn_entries];
+      r = CUDA_CALL_NOCHECK (cuMemcpyDtoH, ind_fn_table, var, bytes);
+      if (r != CUDA_SUCCESS)
+	GOMP_PLUGIN_fatal ("cuMemcpyDtoH error: %s", cuda_error (r));
+
+      /* Build host->target address map for indirect functions.  */
+      uint64_t ind_fn_map[ind_fn_entries * 2 + 1];
+      for (unsigned k = 0; k < ind_fn_entries; k++)
+	{
+	  ind_fn_map[k * 2] = host_ind_fn_table[k];
+	  ind_fn_map[k * 2 + 1] = ind_fn_table[k];
+	  GOMP_PLUGIN_debug (0, "Indirect function %d: %lx->%lx\n",
+			     k, host_ind_fn_table[k], ind_fn_table[k]);
+	}
+      ind_fn_map[ind_fn_entries * 2] = 0;
+
+      /* Write the map onto the target.  */
+      void *map_target_addr
+	= GOMP_OFFLOAD_alloc (ord, sizeof (ind_fn_map));
+      GOMP_PLUGIN_debug (0, "Allocated indirect map at %p\n", map_target_addr);
+
+      GOMP_OFFLOAD_host2dev (ord, map_target_addr,
+			     (void*) ind_fn_map,
+			     sizeof (ind_fn_map));
+
+      /* Write address of the map onto the target.  */
+      CUdeviceptr varptr;
+      size_t varsize;
+      r = CUDA_CALL_NOCHECK (cuModuleGetGlobal, &varptr, &varsize,
+			     module, XSTRING (GOMP_INDIRECT_ADDR_MAP));
+      if (r != CUDA_SUCCESS)
+	GOMP_PLUGIN_fatal ("Indirect map variable not found in image: %s",
+			   cuda_error (r));
+
+      GOMP_PLUGIN_debug (0,
+			 "Indirect map variable found at %llx with size %ld\n",
+			 varptr, varsize);
+
+      GOMP_OFFLOAD_host2dev (ord, (void *) varptr, &map_target_addr,
+			     sizeof (map_target_addr));
+    }
+
   CUdeviceptr varptr;
   size_t varsize;
   CUresult r = CUDA_CALL_NOCHECK (cuModuleGetGlobal, &varptr, &varsize,
diff --git a/libgomp/target.c b/libgomp/target.c
index 812674d19a9..ec154222ad5 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -2256,11 +2256,14 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version,
   void **host_funcs_end  = ((void ***) host_table)[1];
   void **host_var_table  = ((void ***) host_table)[2];
   void **host_vars_end   = ((void ***) host_table)[3];
+  void **host_ind_func_table = ((void ***) host_table)[4];
+  void **host_ind_funcs_end  = ((void ***) host_table)[5];
 
-  /* The func table contains only addresses, the var table contains addresses
-     and corresponding sizes.  */
+  /* The func and ind_func tables contain only addresses, the var table
+     contains addresses and corresponding sizes.  */
   int num_funcs = host_funcs_end - host_func_table;
   int num_vars  = (host_vars_end - host_var_table) / 2;
+  int num_ind_funcs = (host_ind_funcs_end - host_ind_func_table);
 
   /* Load image to device and get target addresses for the image.  */
   struct addr_pair *target_table = NULL;
@@ -2273,7 +2276,9 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version,
   num_target_entries
     = devicep->load_image_func (devicep->target_id, version,
 				target_data, &target_table,
-				rev_lookup ? &rev_target_fn_table : NULL);
+				rev_lookup ? &rev_target_fn_table : NULL,
+				num_ind_funcs
+				  ? (uint64_t *) host_ind_func_table : NULL);
 
   if (num_target_entries != num_funcs + num_vars
       /* "+1" due to the additional ICV struct.  */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare-target-indirect-1.c b/libgomp/testsuite/libgomp.c-c++-common/declare-target-indirect-1.c
new file mode 100644
index 00000000000..b20bfa64dca
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/declare-target-indirect-1.c
@@ -0,0 +1,21 @@
+/* { dg-do run } */
+
+#pragma omp begin declare target indirect
+int foo(void) { return 5; }
+int bar(void) { return 8; }
+int baz(void) { return 11; }
+#pragma omp end declare target
+
+int main (void)
+{
+  int x;
+  int (*foo_ptr) (void) = &foo;
+  int (*bar_ptr) (void) = &bar;
+  int (*baz_ptr) (void) = &baz;
+  int expected = foo () + bar () + baz ();
+
+#pragma omp target map (to: foo_ptr, bar_ptr, baz_ptr) map (from: x)
+  x = (*foo_ptr) () + (*bar_ptr) () + (*baz_ptr) ();
+
+  return x - expected;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare-target-indirect-2.c b/libgomp/testsuite/libgomp.c-c++-common/declare-target-indirect-2.c
new file mode 100644
index 00000000000..9fe190efce8
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/declare-target-indirect-2.c
@@ -0,0 +1,33 @@
+/* { dg-do run } */
+
+#define N 256
+
+#pragma omp begin declare target indirect
+int foo(void) { return 5; }
+int bar(void) { return 8; }
+int baz(void) { return 11; }
+#pragma omp end declare target
+
+int main (void)
+{
+  int i, x = 0, expected = 0;
+  int (*fn_ptr[N])(void);
+
+  for (i = 0; i < N; i++)
+    {
+      switch (i % 3)
+	{
+	case 0: fn_ptr[i] = &foo;
+	case 1: fn_ptr[i] = &bar;
+	case 2: fn_ptr[i] = &baz;
+	}
+      expected += (*fn_ptr[i]) ();
+    }
+
+#pragma omp target teams distribute parallel for reduction(+: x) \
+		map (to: fn_ptr) map (tofrom: x)
+  for (int i = 0; i < N; i++)
+    x += (*fn_ptr[i]) ();
+
+  return x - expected;
+}
-- 
2.34.1


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

end of thread, other threads:[~2024-04-11 10:10 UTC | newest]

Thread overview: 28+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2023-10-08 13:13 [PATCH] openmp: Add support for the 'indirect' clause in C/C++ Kwok Cheung Yeung
2023-10-17 13:12 ` Tobias Burnus
2023-10-17 13:34   ` Jakub Jelinek
2023-10-17 14:41     ` Tobias Burnus
2023-11-03 19:53   ` Kwok Cheung Yeung
2023-11-06  8:48     ` Tobias Burnus
2023-11-07 21:37       ` Joseph Myers
2023-11-07 21:51         ` Jakub Jelinek
2023-11-07 21:59           ` Kwok Cheung Yeung
2023-11-09 12:24     ` Thomas Schwinge
2023-11-09 16:00       ` Tobias Burnus
2023-11-13 10:59         ` Thomas Schwinge
2023-11-13 11:47           ` Tobias Burnus
2024-04-11 10:10             ` Thomas Schwinge
2024-01-03 14:47       ` [committed] " Kwok Cheung Yeung
2024-01-03 15:54       ` Kwok Cheung Yeung
2024-01-22 20:33     ` [PATCH] openmp: Change to using a hashtab to lookup offload target addresses for indirect function calls Kwok Cheung Yeung
2024-01-24  7:06       ` rep.dot.nop
2024-01-29 17:48         ` [PATCH v2] " Kwok Cheung Yeung
2024-03-08 13:40           ` Thomas Schwinge
2024-03-14 11:38           ` Tobias Burnus
2024-01-22 20:41     ` [PATCH] openmp, fortran: Add Fortran support for indirect clause on the declare target directive Kwok Cheung Yeung
2024-01-23 19:14       ` Tobias Burnus
2024-02-05 21:37         ` [PATCH v2] " Kwok Cheung Yeung
2024-02-06  9:03           ` Tobias Burnus
2024-02-06  9:50             ` Kwok Cheung Yeung
2024-02-12  8:51               ` Tobias Burnus
2024-02-15 21:37                 ` [COMMITTED] libgomp: Update documentation for indirect calls in target regions Kwok Cheung Yeung

This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).