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-11-07  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.  Emit error
	message if device_type or indirect clauses used alone.  Emit error
	if indirect clause used with device_type that is not 'any'.
	(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.  Emit error
	message if device_type or indirect clauses used alone.  Emit error
	if indirect clause used with device_type that is not 'any'.
	(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 on an accelerator device.
	* 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-7.c: Update expected error message.
	* c-c++-common/gomp/declare-target-indirect-1.c: New.
	* c-c++-common/gomp/declare-target-indirect-2.c: New.
	* g++.dg/gomp/attrs-21.C (v12): Update expected error message.
	* g++.dg/gomp/declare-target-indirect-1.C: New.
	* gcc.dg/gomp/attrs-21.c (v12): Update expected error message.

include/
	* gomp-constants.h (GOMP_VERSION): Increment to 3.
	(GOMP_VERSION_SUPPORTS_INDIRECT_FUNCS): 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.
	(Other new OpenMP 5.2 features): Add entry for virtual function calls.
	* 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): If the GOMP_VERSION is high
	enough, read host indirect functions table and pass 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.  If the GOMP_VERSION
	is high enough, 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.  If the GOMP_VERSION
	is high enough, 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.
	* testsuite/libgomp.c++/declare-target-indirect-1.C: New.

27 files changed, 551 insertions, 45 deletions

diff --git a/gcc/c-family/c-attribs.cc b/gcc/c-family/c-attribs.cc
index a041c3b..754cdab 100644
--- a/gcc/c-family/c-attribs.cc
+++ b/gcc/c-family/c-attribs.cc
@@ -522,6 +522,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 682157a..9817791 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 4d38750..64d3a94 100644
--- a/gcc/c/c-decl.cc
+++ b/gcc/c/c-decl.cc
@@ -5363,6 +5363,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 09f4d40..0b6db6c 100644
--- a/gcc/c/c-lang.h
+++ b/gcc/c/c-lang.h
@@ -63,6 +63,7 @@ struct GTY(()) language_function {
 struct GTY(()) c_omp_declare_target_attr {
   bool attr_syntax;
   int device_type;
+  int indirect;
 };
 
 struct GTY(()) c_omp_begin_assumes_data {
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 134d3ed..703f957 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -14598,6 +14598,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))
@@ -15474,6 +15476,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 logical 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 )
 
@@ -19035,6 +19078,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";
@@ -24608,14 +24655,16 @@ c_maybe_parse_omp_decl (tree decl, tree d)
 	( (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 only_device_type = true;
+  bool indirect = false;
+  bool only_device_type_or_indirect = true;
   if (c_parser_next_token_is (parser, CPP_NAME)
       || (c_parser_next_token_is (parser, CPP_COMMA)
 	  && c_parser_peek_2nd_token (parser)->type == CPP_NAME))
@@ -24633,22 +24682,27 @@ c_parser_omp_declare_target (c_parser *parser)
     {
       bool attr_syntax = parser->in_omp_attribute_pragma != NULL;
       c_parser_skip_to_pragma_eol (parser);
-      c_omp_declare_target_attr attr = { attr_syntax, -1 };
+      c_omp_declare_target_attr attr = { attr_syntax, -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));
       tree at2 = lookup_attribute ("omp declare target link",
 				   DECL_ATTRIBUTES (t));
-      only_device_type = false;
+      only_device_type_or_indirect = false;
       if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LINK)
 	{
 	  id = get_identifier ("omp declare target link");
@@ -24710,10 +24764,25 @@ 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)
+  if ((device_type || indirect) && only_device_type_or_indirect)
     error_at (OMP_CLAUSE_LOCATION (clauses),
-	      "directive with only %<device_type%> clause");
+	      "directive with only %<device_type%> or %<indirect%> clauses");
+  if (indirect && device_type && device_type != OMP_CLAUSE_DEVICE_TYPE_ANY)
+    error_at (OMP_CLAUSE_LOCATION (clauses),
+	      "%<device_type%> clause must specify 'any' when used with "
+	      "an %<indirect%> clause");
 }
 
 /* OpenMP 5.1
@@ -24722,7 +24791,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)
@@ -24746,10 +24816,16 @@ 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 = { attr_syntax, 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 = { attr_syntax, 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 9c87945..4580ff0 100644
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -15914,6 +15914,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 f6d56b7..0e224ca 100644
--- a/gcc/config/gcn/mkoffload.cc
+++ b/gcc/config/gcn/mkoffload.cc
@@ -479,7 +479,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);
@@ -508,7 +509,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))
     {
@@ -570,6 +572,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;
@@ -597,6 +610,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;
@@ -634,6 +656,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"
@@ -728,12 +751,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 aaea9fb..fb75ca0 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 634c316..0eeff95 100644
--- a/gcc/config/nvptx/nvptx.cc
+++ b/gcc/config/nvptx/nvptx.cc
@@ -5919,7 +5919,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 98b29e9..b2603d4 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 0aa1e35..9e666e5 100644
--- a/gcc/cp/decl2.cc
+++ b/gcc/cp/decl2.cc
@@ -1762,6 +1762,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 0fff981..5116bcb 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -37524,6 +37524,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))
@@ -38558,6 +38560,46 @@ 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;
+
+      bool non_constant_p;
+      t = cp_parser_constant_expression (parser, true, &non_constant_p);
+
+      if (t != error_mark_node && non_constant_p)
+	error_at (location, "expected constant logical 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 )
 
@@ -41629,6 +41671,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";
@@ -48171,7 +48218,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));
@@ -48235,6 +48283,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;
 }
 
@@ -48252,14 +48311,16 @@ 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 only_device_type = true;
+  bool indirect = false;
+  bool only_device_type_or_indirect = true;
   if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)
       || (cp_lexer_next_token_is (parser->lexer, CPP_COMMA)
 	  && cp_lexer_nth_token_is (parser->lexer, 2, CPP_NAME)))
@@ -48277,21 +48338,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))
+      only_device_type_or_indirect = false;
+      if (!handle_omp_declare_target_clause (c, t, device_type, indirect))
 	continue;
       if (VAR_OR_FUNCTION_DECL_P (t)
 	  && DECL_LOCAL_DECL_P (t)
@@ -48299,11 +48365,15 @@ 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)
+  if ((device_type || indirect) && only_device_type_or_indirect)
+    error_at (OMP_CLAUSE_LOCATION (clauses),
+	      "directive with only %<device_type%> or %<indirect%> clauses");
+  if (indirect && device_type && device_type != OMP_CLAUSE_DEVICE_TYPE_ANY)
     error_at (OMP_CLAUSE_LOCATION (clauses),
-	      "directive with only %<device_type%> clause");
+	      "%<device_type%> clause must specify 'any' when used with "
+	      "an %<indirect%> clause");
 }
 
 /* OpenMP 5.1
@@ -48312,7 +48382,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)
@@ -48342,11 +48413,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 37bffca..4059e74 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -8888,6 +8888,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 32c0f5a..db6a22a 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 aa1b2f2..f7ed622 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 e0f0326..ed78d49 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 0d3c879..1d6dfef 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,11 @@ execute_omp_device_lower ()
   gimple_stmt_iterator gsi;
   bool calls_declare_variant_alt
     = cgraph_node::get (cfun->decl)->calls_declare_variant_alt;
+#ifdef ACCEL_COMPILER
+  bool omp_redirect_indirect_calls = vec_safe_length (offload_ind_funcs) > 0;
+  tree map_ptr_fn
+    = builtin_decl_explicit (BUILT_IN_GOMP_TARGET_MAP_INDIRECT_PTR);
+#endif
   FOR_EACH_BB_FN (bb, cfun)
     for (gsi = gsi_start_bb (bb); !gsi_end_p (gsi); gsi_next (&gsi))
       {
@@ -2621,6 +2655,33 @@ execute_omp_device_lower ()
 		      update_stmt (stmt);
 		    }
 		}
+#ifdef ACCEL_COMPILER
+	    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 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);
+	      }
+#endif
 	    continue;
 	  }
 	tree lhs = gimple_call_lhs (stmt), rhs = NULL_TREE;
@@ -2757,9 +2818,15 @@ public:
   /* opt_pass methods: */
   bool gate (function *fun) final override
     {
+#ifdef ACCEL_COMPILER
+      bool offload_ind_funcs_p = vec_safe_length (offload_ind_funcs) > 0;
+#else
+      bool offload_ind_funcs_p = false;
+#endif
       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
+		      || offload_ind_funcs_p)));
     }
   unsigned int execute (function *) final override
     {
diff --git a/gcc/omp-offload.h b/gcc/omp-offload.h
index 73711e7..ae36442 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-7.c b/gcc/testsuite/c-c++-common/gomp/declare-target-7.c
index 747000a..e37b465 100644
--- a/gcc/testsuite/c-c++-common/gomp/declare-target-7.c
+++ b/gcc/testsuite/c-c++-common/gomp/declare-target-7.c
@@ -1,7 +1,7 @@
 /* { dg-do compile } */
 /* { dg-options "-fopenmp" } */
 
-#pragma omp declare target device_type (any)		/* { dg-error "directive with only 'device_type' clause" } */
+#pragma omp declare target device_type (any)		/* { dg-error "directive with only 'device_type' or 'indirect' clauses" } */
 
 void f1 (void) {}
 #pragma omp declare target device_type (host) to (f1) device_type (nohost)	/* { dg-error "too many 'device_type' clauses" } */
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 0000000..0fcbb2d
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/declare-target-indirect-1.c
@@ -0,0 +1,62 @@
+/* { 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 logical 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 /* { dg-error "directive with only 'device_type' or 'indirect' clauses" }*/
+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" } */
+
+void fn12 (void) { }
+#pragma omp  declare target indirect ("abs") to (fn12)
+
+void fn13 (void) { }
+#pragma omp declare target indirect (5.5) enter (fn13)
+
+void fn14 (void) { }
+#pragma omp declare target indirect (1) device_type (host) enter (fn14) /* { dg-error "'device_type' clause must specify 'any' when used with an 'indirect' clause" } */
+
+void fn15 (void) { }
+#pragma omp declare target indirect (0) device_type (nohost) enter (fn15)
+
+/* 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 0000000..6ba278b
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/declare-target-indirect-2.c
@@ -0,0 +1,32 @@
+/* { 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" } } */
+
+#pragma omp begin declare target indirect(1)
+  int foo(void) { return 5; }
+  #pragma omp begin declare target indirect(0)
+    int bar(void) { return 8; }
+    int baz(void) { return 11; }
+    #pragma omp declare target indirect enter(baz)
+  #pragma omp end declare target
+#pragma omp end declare target
+/* { dg-final { scan-tree-dump "__attribute__\\\(\\\(omp declare target, omp declare target block, omp declare target indirect\\\)\\\)\\\nint foo" "gimple" } } */
+/* { dg-final { scan-tree-dump "__attribute__\\\(\\\(omp declare target, omp declare target block\\\)\\\)\\\nint bar" "gimple" } } */
+/* { dg-final { scan-tree-dump "__attribute__\\\(\\\(omp declare target indirect, omp declare target, omp declare target block\\\)\\\)\\\nint baz" "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/attrs-21.C b/gcc/testsuite/g++.dg/gomp/attrs-21.C
index 46bdef2..03c760f 100644
--- a/gcc/testsuite/g++.dg/gomp/attrs-21.C
+++ b/gcc/testsuite/g++.dg/gomp/attrs-21.C
@@ -20,7 +20,7 @@ foo ()
   [[omp::decl (declare target (v8))]] static int v9;		// { dg-error "expected end of line before '\\\(' token" }
   [[omp::decl (declare target enter (v8))]] static int v10;	// { dg-error "expected an OpenMP clause before '\\\(' token" }
   [[omp::decl (declare target, link (v9))]] static int v11;	// { dg-error "expected an OpenMP clause before '\\\(' token" }
-  [[omp::decl (declare target device_type (any))]] static int v12;	// { dg-error "directive with only 'device_type' clause" }
+  [[omp::decl (declare target device_type (any))]] static int v12;	// { dg-error "directive with only 'device_type' or 'indirect' clauses" }
 }
 
 int i;
diff --git a/gcc/testsuite/g++.dg/gomp/declare-target-indirect-1.C b/gcc/testsuite/g++.dg/gomp/declare-target-indirect-1.C
new file mode 100644
index 0000000..1d66ec9
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/declare-target-indirect-1.C
@@ -0,0 +1,17 @@
+// { dg-skip-if "c++98 does not support attributes" { c++98_only } }
+
+[[omp::decl (declare target, indirect(1))]] // { dg-error "directive with only 'device_type' or 'indirect' clause" }
+int f (void) { return 5; }
+
+[[omp::decl (declare target indirect)]] // { dg-error "directive with only 'device_type' or 'indirect' clause" }
+int g (void) { return 8; }
+
+[[omp::directive (begin declare target, indirect)]];
+int h (void) { return 11; }
+[[omp::directive (end declare target)]];
+
+int i (void) { return 8; }
+[[omp::directive (declare target to(i), indirect (1))]];
+
+int j (void) { return 11; }
+[[omp::directive (declare target indirect enter (j))]];
diff --git a/gcc/testsuite/gcc.dg/gomp/attrs-21.c b/gcc/testsuite/gcc.dg/gomp/attrs-21.c
index bd8ff11..ca97b76 100644
--- a/gcc/testsuite/gcc.dg/gomp/attrs-21.c
+++ b/gcc/testsuite/gcc.dg/gomp/attrs-21.c
@@ -21,7 +21,7 @@ foo ()
   [[omp::decl (declare target (v8))]] static int v9;		/* { dg-error "expected end of line before '\\\(' token" } */
   [[omp::decl (declare target enter (v8))]] static int v10;	/* { dg-error "expected an OpenMP clause before '\\\(' token" } */
   [[omp::decl (declare target, link (v9))]] static int v11;	/* { dg-error "expected an OpenMP clause before '\\\(' token" } */
-  [[omp::decl (declare target device_type (any))]] static int v12;	/* { dg-error "directive with only 'device_type' clause" } */
+  [[omp::decl (declare target device_type (any))]] static int v12;	/* { dg-error "directive with only 'device_type' or 'indirect' clauses" } */
 }
 
 int i;
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index 1343534..65e51b9 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -350,6 +350,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 9c9b057..33ea1d2 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  */
@@ -361,6 +362,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 632a1af..086b55f 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)