OpenMP: Add uses_allocators support

2023-11-19  Tobias Burnus  <tobias@codesourcery.com>
	    Chung-Lin Tang <cltang@codesourcery.com>

gcc/ChangeLog:

	* builtin-types.def (BT_FN_VOID_PTRMODE):
	(BT_FN_PTRMODE_PTRMODE_INT_PTR): Add.
	* gimplify.cc (gimplify_bind_expr): Diagnose missing
	uses_allocators clause.
	(gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses,
	gimplify_omp_workshare): Handle uses_allocators.
	* omp-builtins.def (BUILT_IN_OMP_INIT_ALLOCATOR,
	BUILT_IN_OMP_DESTROY_ALLOCATOR): Add.
	* omp-low.cc (scan_sharing_clauses):
	* tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_USES_ALLOCATORS.
	* tree.cc (omp_clause_num_ops, omp_clause_code_name): Likewise.
	* tree-pretty-print.cc (dump_omp_clause): Handle it.
	* tree.h (OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR,
	OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE,
	OMP_CLAUSE_USES_ALLOCATORS_TRAITS): New.

gcc/c-family/ChangeLog:

	* c-omp.cc (c_omp_split_clauses): Hande uses_allocators.
	* c-pragma.h (enum pragma_omp_clause): Add
	PRAGMA_OMP_CLAUSE_USES_ALLOCATORS.

gcc/c/ChangeLog:

	* c-parser.cc (c_parser_omp_clause_uses_allocators): New.
	(c_parser_omp_clause_name, c_parser_omp_all_clauses,
	OMP_TARGET_CLAUSE_MASK): Handle uses_allocators.
	* c-typeck.cc (c_finish_omp_clauses): Likewise.

gcc/cp/ChangeLog:

	* parser.cc (cp_parser_omp_clause_uses_allocators): New.
	(cp_parser_omp_clause_name, cp_parser_omp_all_clauses,
	OMP_TARGET_CLAUSE_MASK): Handle uses_allocators.
	* semantics.cc (finish_omp_clauses): Likewise.

gcc/fortran/ChangeLog:

	* trans-array.cc (gfc_conv_array_initializer): Set PURPOSE
	when building constructor for get_initialized_tmp_var.
	* trans-openmp.cc (gfc_trans_omp_clauses): Handle uses_allocators.
	* types.def (BT_FN_VOID_PTRMODE, BT_FN_PTRMODE_PTRMODE_INT_PTR): Add.

libgomp/ChangeLog:

	* testsuite/libgomp.c++/c++.exp (check_effective_target_c,
	check_effective_target_c++): Add.
	* testsuite/libgomp.c/c.exp (check_effective_target_c,
	check_effective_target_c++): Add.
	* testsuite/libgomp.fortran/uses_allocators_2.f90: Remove 'sorry'.
	* testsuite/libgomp.c-c++-common/uses_allocators-1.c: New test.
	* testsuite/libgomp.c-c++-common/uses_allocators-2.c: New test.
	* testsuite/libgomp.c-c++-common/uses_allocators-3.c: New test.
	* testsuite/libgomp.c-c++-common/uses_allocators-4.c: New test.
	* testsuite/libgomp.fortran/uses_allocators_3.f90: New test.
	* testsuite/libgomp.fortran/uses_allocators_4.f90: New test.
	* testsuite/libgomp.fortran/uses_allocators_5.f90: New test.
	* testsuite/libgomp.fortran/uses_allocators_6.f90: New test.

gcc/testsuite/ChangeLog:

	* gfortran.dg/gomp/allocate-1.f90: Add uses_allocators.
	* gfortran.dg/gomp/scope-6.f90: Update dg-scan-tree-dump.
	* c-c++-common/gomp/uses_allocators-1.c: New test.
	* c-c++-common/gomp/uses_allocators-2.c: New test.
	* gfortran.dg/gomp/uses_allocators-1.f90: New test.

Co-authored-by: Chung-Lin Tang <cltang@codesourcery.com>

 gcc/builtin-types.def                              |   3 +
 gcc/c-family/c-omp.cc                              |   1 +
 gcc/c-family/c-pragma.h                            |   1 +
 gcc/c/c-parser.cc                                  | 216 ++++++++++++++++++-
 gcc/c/c-typeck.cc                                  | 105 +++++++++
 gcc/cp/parser.cc                                   | 237 ++++++++++++++++++++-
 gcc/cp/semantics.cc                                |  95 +++++++++
 gcc/fortran/trans-array.cc                         |   5 +-
 gcc/fortran/trans-openmp.cc                        |  42 +++-
 gcc/fortran/types.def                              |   3 +
 gcc/gimplify.cc                                    | 183 +++++++++++++++-
 gcc/omp-builtins.def                               |   4 +
 gcc/omp-low.cc                                     |  32 +++
 .../c-c++-common/gomp/uses_allocators-1.c          |  46 ++++
 .../c-c++-common/gomp/uses_allocators-2.c          |  33 +++
 gcc/testsuite/gfortran.dg/gomp/allocate-1.f90      |   7 +-
 gcc/testsuite/gfortran.dg/gomp/scope-6.f90         |   2 +-
 .../gfortran.dg/gomp/uses_allocators-1.f90         |  23 ++
 gcc/tree-core.h                                    |   3 +
 gcc/tree-pretty-print.cc                           |  14 ++
 gcc/tree.cc                                        |   2 +
 gcc/tree.h                                         |   9 +
 libgomp/testsuite/libgomp.c++/c++.exp              |   9 +
 .../libgomp.c-c++-common/uses_allocators-1.c       |  53 +++++
 .../libgomp.c-c++-common/uses_allocators-2.c       |  39 ++++
 .../libgomp.c-c++-common/uses_allocators-3.c       |  37 ++++
 .../libgomp.c-c++-common/uses_allocators-4.c       |  53 +++++
 libgomp/testsuite/libgomp.c/c.exp                  |   8 +
 .../libgomp.fortran/uses_allocators_2.f90          |  22 +-
 .../libgomp.fortran/uses_allocators_3.f90          |  62 ++++++
 .../libgomp.fortran/uses_allocators_4.f90          |  54 +++++
 .../libgomp.fortran/uses_allocators_5.f90          |  14 ++
 .../libgomp.fortran/uses_allocators_6.f90          |  50 +++++
 33 files changed, 1435 insertions(+), 32 deletions(-)

diff --git a/gcc/builtin-types.def b/gcc/builtin-types.def
index 43381bc8949..da9719272ec 100644
--- a/gcc/builtin-types.def
+++ b/gcc/builtin-types.def
@@ -385,6 +385,7 @@ DEF_FUNCTION_TYPE_1 (BT_FN_DFLOAT32_DFLOAT32, BT_DFLOAT32, BT_DFLOAT32)
 DEF_FUNCTION_TYPE_1 (BT_FN_DFLOAT64_DFLOAT64, BT_DFLOAT64, BT_DFLOAT64)
 DEF_FUNCTION_TYPE_1 (BT_FN_DFLOAT128_DFLOAT128, BT_DFLOAT128, BT_DFLOAT128)
 DEF_FUNCTION_TYPE_1 (BT_FN_VOID_VPTR, BT_VOID, BT_VOLATILE_PTR)
+DEF_FUNCTION_TYPE_1 (BT_FN_VOID_PTRMODE, BT_VOID, BT_PTRMODE)
 DEF_FUNCTION_TYPE_1 (BT_FN_VOID_PTRPTR, BT_VOID, BT_PTR_PTR)
 DEF_FUNCTION_TYPE_1 (BT_FN_VOID_CONST_PTR, BT_VOID, BT_CONST_PTR)
 DEF_FUNCTION_TYPE_1 (BT_FN_UINT_UINT, BT_UINT, BT_UINT)
@@ -829,6 +830,8 @@ DEF_FUNCTION_TYPE_3 (BT_FN_PTR_SIZE_SIZE_PTRMODE,
 		     BT_PTR, BT_SIZE, BT_SIZE, BT_PTRMODE)
 DEF_FUNCTION_TYPE_3 (BT_FN_VOID_PTR_UINT8_PTRMODE, BT_VOID, BT_PTR, BT_UINT8,
 		     BT_PTRMODE)
+DEF_FUNCTION_TYPE_3 (BT_FN_PTRMODE_PTRMODE_INT_PTR, BT_PTRMODE, BT_PTRMODE,
+		     BT_INT, BT_PTR)
 
 DEF_FUNCTION_TYPE_4 (BT_FN_SIZE_CONST_PTR_SIZE_SIZE_FILEPTR,
 		     BT_SIZE, BT_CONST_PTR, BT_SIZE, BT_SIZE, BT_FILEPTR)
diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc
index 95b6c1e623f..82422fc8c3e 100644
--- a/gcc/c-family/c-omp.cc
+++ b/gcc/c-family/c-omp.cc
@@ -2026,6 +2026,7 @@ c_omp_split_clauses (location_t loc, enum tree_code code,
 	case OMP_CLAUSE_HAS_DEVICE_ADDR:
 	case OMP_CLAUSE_DEFAULTMAP:
 	case OMP_CLAUSE_DEPEND:
+	case OMP_CLAUSE_USES_ALLOCATORS:
 	  s = C_OMP_CLAUSE_SPLIT_TARGET;
 	  break;
 	case OMP_CLAUSE_DOACROSS:
diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index 98177913053..7b1d9dc7232 100644
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -161,6 +161,7 @@ enum pragma_omp_clause {
   PRAGMA_OMP_CLAUSE_UNTIED,
   PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR,
   PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR,
+  PRAGMA_OMP_CLAUSE_USES_ALLOCATORS,
 
   /* Clauses for OpenACC.  */
   PRAGMA_OACC_CLAUSE_ASYNC,
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 703f9570dbc..db69a418f59 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -14722,6 +14722,8 @@ c_parser_omp_clause_name (c_parser *parser)
 	    result = PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR;
 	  else if (!strcmp ("use_device_ptr", p))
 	    result = PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR;
+	  else if (!strcmp ("uses_allocators", p))
+	    result = PRAGMA_OMP_CLAUSE_USES_ALLOCATORS;
 	  break;
 	case 'v':
 	  if (!strcmp ("vector", p))
@@ -17568,6 +17570,213 @@ c_parser_omp_clause_allocate (c_parser *parser, tree list)
   return nl;
 }
 
+/* OpenMP 5.0:
+   uses_allocators ( allocator-list )
+
+   allocator-list:
+   allocator
+   allocator , allocator-list
+   allocator ( traits-array )
+   allocator ( traits-array ) , allocator-list
+
+   OpenMP 5.2:
+
+   uses_allocators ( modifier : allocator-list )
+   uses_allocators ( modifier , modifier : allocator-list )
+
+   modifier:
+   traits ( traits-array )
+   memspace ( mem-space-handle )  */
+
+static tree
+c_parser_omp_clause_uses_allocators (c_parser *parser, tree list)
+{
+  location_t clause_loc = c_parser_peek_token (parser)->location;
+  tree t = NULL_TREE, nl = list;
+  matching_parens parens;
+  if (!parens.require_open (parser))
+    return list;
+
+  tree memspace_expr = NULL_TREE;
+  tree traits_var = NULL_TREE;
+
+  struct item_tok
+  {
+    location_t loc;
+    tree id;
+    item_tok (void) : loc (UNKNOWN_LOCATION), id (NULL_TREE) {}
+  };
+  struct item { item_tok name, arg; };
+  auto_vec<item> *modifiers = NULL, *allocators = NULL;
+  auto_vec<item> *cur_list = new auto_vec<item> (4);
+
+  while (true)
+    {
+      item it;
+
+      if (c_parser_next_token_is (parser, CPP_NAME))
+	{
+	  c_token *tok = c_parser_peek_token (parser);
+	  it.name.id = tok->value;
+	  it.name.loc = tok->location;
+	  c_parser_consume_token (parser);
+
+	  if (c_parser_next_token_is (parser, CPP_OPEN_PAREN))
+	    {
+	      matching_parens parens2;
+	      parens2.consume_open (parser);
+
+	      if (c_parser_next_token_is (parser, CPP_NAME))
+		{
+		  tok = c_parser_peek_token (parser);
+		  it.arg.id = tok->value;
+		  it.arg.loc = tok->location;
+		  c_parser_consume_token (parser);
+		}
+	      else
+		{
+		  c_parser_error (parser, "expected identifier");
+		  parens2.skip_until_found_close (parser);
+		  goto end;
+		}
+	      parens2.skip_until_found_close (parser);
+	    }
+	}
+
+      cur_list->safe_push (it);
+
+      if (c_parser_next_token_is (parser, CPP_COMMA))
+	c_parser_consume_token (parser);
+      else if (c_parser_next_token_is (parser, CPP_COLON))
+	{
+	  if (modifiers)
+	    {
+	      c_parser_error (parser, "expected %<)%>");
+	      goto end;
+	    }
+	  else
+	    {
+	      c_parser_consume_token (parser);
+	      modifiers = cur_list;
+	      cur_list = new auto_vec<item> (4);
+	    }
+	}
+      else if (c_parser_next_token_is (parser, CPP_CLOSE_PAREN))
+	{
+	  gcc_assert (allocators == NULL);
+	  allocators = cur_list;
+	  cur_list = NULL;
+	  break;
+	}
+      else
+	{
+	  c_parser_error (parser, "expected %<)%>");
+	  goto end;
+	}
+    }
+
+  if (modifiers)
+    for (unsigned i = 0; i < modifiers->length (); i++)
+      {
+	item& it = (*modifiers)[i];
+	const char *p = IDENTIFIER_POINTER (it.name.id);
+	int strcmp_traits = 1, strcmp_memspace = 1;
+
+	if ((strcmp_traits = strcmp ("traits", p)) == 0
+	    || (strcmp_memspace = strcmp ("memspace", p)) == 0)
+	  {
+	    if ((strcmp_traits == 0 && traits_var != NULL_TREE)
+		|| (strcmp_memspace == 0 && memspace_expr != NULL_TREE))
+	      {
+		error_at (it.name.loc, "duplicate %qs modifier", p);
+		goto end;
+	      }
+	    t = lookup_name (it.arg.id);
+	    if (t == NULL_TREE)
+	      {
+		undeclared_variable (it.arg.loc, it.arg.id);
+		t = error_mark_node;
+	      }
+	    else if (strcmp_memspace == 0)
+	      memspace_expr = t;
+	    else if (strcmp_traits == 0)
+	      traits_var = t;
+	    else
+	      gcc_unreachable ();
+	  }
+	else
+	  {
+	    error_at (it.name.loc, "unknown modifier %qE", it.name.id);
+	    goto end;
+	  }
+      }
+
+  if (allocators)
+    {
+      if (modifiers)
+	{
+	  if (allocators->length () > 1)
+	    {
+	      error_at ((*allocators)[1].name.loc,
+			"%<uses_allocators%> clause only accepts a single "
+			"allocator when using modifiers");
+	      goto end;
+	    }
+	  else if ((*allocators)[0].arg.id)
+	    {
+	      error_at ((*allocators)[0].arg.loc,
+			"legacy %<%E(%E)%> traits syntax not allowed in "
+			"%<uses_allocators%> clause when using modifiers",
+			(*allocators)[0].name.id, (*allocators)[0].arg.id);
+	      goto end;
+	    }
+	}
+
+      for (unsigned i = 0; i < allocators->length (); i++)
+	{
+	  item& it = (*allocators)[i];
+	  t = lookup_name (it.name.id);
+	  if (t == NULL_TREE)
+	    {
+	      undeclared_variable (it.name.loc, it.name.id);
+	      goto end;
+	    }
+	  else if (t != error_mark_node)
+	    {
+	      tree t2 = NULL_TREE;
+	      if (it.arg.id)
+		{
+		  t2 = lookup_name (it.arg.id);
+		  if (t2 == NULL_TREE)
+		    {
+		      undeclared_variable (it.arg.loc, it.arg.id);
+		      goto end;
+		    }
+		}
+	      else
+		t2 = traits_var;
+
+	      tree c = build_omp_clause (clause_loc,
+					 OMP_CLAUSE_USES_ALLOCATORS);
+	      OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c) = t;
+	      OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c) = memspace_expr;
+	      OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c) = t2;
+	      OMP_CLAUSE_CHAIN (c) = nl;
+	      nl = c;
+	    }
+	}
+    }
+ end:
+  if (cur_list)
+    delete cur_list;
+  if (modifiers)
+    delete modifiers;
+  if (allocators)
+    delete allocators;
+  parens.skip_until_found_close (parser);
+  return nl;
+}
+
 /* OpenMP 4.0:
    linear ( variable-list )
    linear ( variable-list : expression )
@@ -19237,6 +19446,10 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask,
 	  clauses = c_parser_omp_clause_allocate (parser, clauses);
 	  c_name = "allocate";
 	  break;
+	case PRAGMA_OMP_CLAUSE_USES_ALLOCATORS:
+	  clauses = c_parser_omp_clause_uses_allocators (parser, clauses);
+	  c_name = "uses_allocators";
+	  break;
 	case PRAGMA_OMP_CLAUSE_LINEAR: 
 	  clauses = c_parser_omp_clause_linear (parser, clauses); 
 	  c_name = "linear";
@@ -23648,7 +23861,8 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser,
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IN_REDUCTION)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_THREAD_LIMIT)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR)\
-	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR))
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR)\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USES_ALLOCATORS))
 
 static bool
 c_parser_omp_target (c_parser *parser, enum pragma_context context, bool *if_p)
diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc
index 1dbb4471a88..e17414a5245 100644
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -15136,6 +15136,111 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	      break;
 	    }
 	  gcc_unreachable ();
+
+	case OMP_CLAUSE_USES_ALLOCATORS:
+	  t = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c);
+	  if ((VAR_P (t) || TREE_CODE (t) == PARM_DECL)
+	      && (bitmap_bit_p (&generic_head, DECL_UID (t))
+		  || bitmap_bit_p (&map_head, DECL_UID (t))
+		  || bitmap_bit_p (&firstprivate_head, DECL_UID (t))
+		  || bitmap_bit_p (&lastprivate_head, DECL_UID (t))))
+	    {
+	      error_at (OMP_CLAUSE_LOCATION (c),
+			"%qE appears more than once in data clauses", t);
+	      remove = true;
+	      break;
+	    }
+	  else
+	    bitmap_set_bit (&generic_head, DECL_UID (t));
+	  if (TREE_CODE (TREE_TYPE (t)) != ENUMERAL_TYPE
+	      || strcmp (IDENTIFIER_POINTER (TYPE_IDENTIFIER (TREE_TYPE (t))),
+			 "omp_allocator_handle_t") != 0)
+	    {
+	      error_at (OMP_CLAUSE_LOCATION (c),
+			"allocator must be of %<omp_allocator_handle_t%> type");
+	      remove = true;
+	      break;
+	    }
+	  if (TREE_CODE (t) == CONST_DECL)
+	    {
+	      /* Currently for pre-defined allocators in libgomp, we do not
+		 require additional init/fini inside target regions, so discard
+		 such clauses.  */
+	      remove = true;
+
+	      if (strcmp (IDENTIFIER_POINTER (DECL_NAME (t)),
+			  "omp_null_allocator") == 0)
+		{
+		  error_at (OMP_CLAUSE_LOCATION (c),
+			    "%<omp_null_allocator%> cannot be used in "
+			    "%<uses_allocators%> clause");
+		  break;
+		}
+
+	      if (OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c)
+		  || OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c))
+		{
+		  error_at (OMP_CLAUSE_LOCATION (c),
+			    "modifiers cannot be used with pre-defined "
+			    "allocators");
+		  break;
+		}
+	    }
+	  t = OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c);
+	  if (t != NULL_TREE
+	      && (TREE_CODE (t) != CONST_DECL
+		  || TREE_CODE (TREE_TYPE (t)) != ENUMERAL_TYPE
+		  || strcmp (IDENTIFIER_POINTER (TYPE_IDENTIFIER (TREE_TYPE (t))),
+			     "omp_memspace_handle_t") != 0))
+	    {
+	      error_at (OMP_CLAUSE_LOCATION (c), "memspace modifier must be "
+			"constant enum of %<omp_memspace_handle_t%> type");
+	      remove = true;
+	      break;
+	    }
+	  t = OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c);
+	  if (t != NULL_TREE)
+	    {
+	      bool type_err = false;
+
+	      if (TREE_CODE (TREE_TYPE (t)) != ARRAY_TYPE
+		  || DECL_SIZE (t) == NULL_TREE)
+		type_err = true;
+	      else
+		{
+		  tree elem_t = TREE_TYPE (TREE_TYPE (t));
+		  if (TREE_CODE (elem_t) != RECORD_TYPE
+		      || strcmp (IDENTIFIER_POINTER (TYPE_IDENTIFIER (elem_t)),
+				 "omp_alloctrait_t") != 0
+		      || !TYPE_READONLY (elem_t))
+		    type_err = true;
+		}
+	      if (type_err)
+		{
+		  if (TREE_CODE (t) != ERROR_MARK)
+		    error_at (OMP_CLAUSE_LOCATION (c), "traits array %qE must "
+			      "be of %<const omp_alloctrait_t []%> type", t);
+		  else
+		    error_at (OMP_CLAUSE_LOCATION (c), "traits array must "
+			      "be of %<const omp_alloctrait_t []%> type");
+		  remove = true;
+		}
+	      else
+		{
+		  tree cst_val = decl_constant_value_1 (t, true);
+		  if (cst_val == t)
+		    {
+		      error_at (OMP_CLAUSE_LOCATION (c), "traits array must be "
+				"of constant values");
+
+		      remove = true;
+		    }
+		}
+	    }
+	  if (remove)
+	    break;
+	  pc = &OMP_CLAUSE_CHAIN (c);
+	  continue;
 	case OMP_CLAUSE_DEPEND:
 	case OMP_CLAUSE_AFFINITY:
 	  t = OMP_CLAUSE_DECL (c);
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index f556b8f3c01..15216a08b54 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -37657,6 +37657,8 @@ cp_parser_omp_clause_name (cp_parser *parser)
 	    result = PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR;
 	  else if (!strcmp ("use_device_ptr", p))
 	    result = PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR;
+	  else if (!strcmp ("uses_allocators", p))
+	    result = PRAGMA_OMP_CLAUSE_USES_ALLOCATORS;
 	  break;
 	case 'v':
 	  if (!strcmp ("vector", p))
@@ -39983,6 +39985,234 @@ cp_parser_omp_clause_allocate (cp_parser *parser, tree list)
   return nlist;
 }
 
+/* OpenMP 5.0:
+   uses_allocators ( allocator-list )
+
+   allocator-list:
+   allocator
+   allocator , allocator-list
+   allocator ( traits-array )
+   allocator ( traits-array ) , allocator-list
+
+   OpenMP 5.2:
+
+   uses_allocators ( modifier : allocator-list )
+   uses_allocators ( modifier , modifier : allocator-list )
+
+   modifier:
+   traits ( traits-array )
+   memspace ( mem-space-handle )  */
+
+static tree
+cp_parser_omp_clause_uses_allocators (cp_parser *parser, tree list)
+{
+  location_t clause_loc
+    = cp_lexer_peek_token (parser->lexer)->location;
+  tree t = NULL_TREE, nl = list;
+  matching_parens parens;
+  if (!parens.require_open (parser))
+    return list;
+
+  tree memspace_expr = NULL_TREE;
+  tree traits_var = NULL_TREE;
+
+  struct item_tok
+  {
+    location_t loc;
+    tree id;
+    item_tok (void) : loc (UNKNOWN_LOCATION), id (NULL_TREE) {}
+  };
+  struct item { item_tok name, arg; };
+  auto_vec<item> *modifiers = NULL, *allocators = NULL;
+  auto_vec<item> *cur_list = new auto_vec<item> (4);
+
+  while (true)
+    {
+      item it;
+
+      if (cp_lexer_next_token_is (parser->lexer, CPP_NAME))
+	{
+	  cp_token *tok = cp_lexer_peek_token (parser->lexer);
+	  it.name.id = tok->u.value;
+	  it.name.loc = tok->location;
+	  cp_lexer_consume_token (parser->lexer);
+
+	  if (cp_lexer_next_token_is (parser->lexer, CPP_OPEN_PAREN))
+	    {
+	      matching_parens parens2;
+	      parens2.consume_open (parser);
+
+	      if (cp_lexer_next_token_is (parser->lexer, CPP_NAME))
+		{
+		  tok = cp_lexer_peek_token (parser->lexer);
+		  it.arg.id = tok->u.value;
+		  it.arg.loc = tok->location;
+		  cp_lexer_consume_token (parser->lexer);
+		}
+	      else
+		{
+		  cp_parser_error (parser, "expected identifier");
+		  cp_parser_skip_to_closing_parenthesis (parser,
+							 /*recovering=*/true,
+							 /*or_comma=*/false,
+							 /*consume_paren=*/true);
+		  goto end;
+		}
+	      cp_parser_skip_to_closing_parenthesis (parser,
+						     /*recovering=*/false,
+						     /*or_comma=*/false,
+						     /*consume_paren=*/true);
+	    }
+	}
+
+      cur_list->safe_push (it);
+
+      if (cp_lexer_next_token_is (parser->lexer, CPP_COMMA))
+	cp_lexer_consume_token (parser->lexer);
+      else if (cp_lexer_next_token_is (parser->lexer, CPP_COLON))
+	{
+	  if (modifiers)
+	    {
+	      cp_parser_error (parser, "expected %<)%>");
+	      cp_parser_skip_to_closing_parenthesis (parser,
+						     /*recovering=*/true,
+						     /*or_comma=*/false,
+						     /*consume_paren=*/true);
+	      goto end;
+	    }
+	  else
+	    {
+	      cp_lexer_consume_token (parser->lexer);
+	      modifiers = cur_list;
+	      cur_list = new auto_vec<item> (4);
+	    }
+	}
+      else if (cp_lexer_next_token_is (parser->lexer, CPP_CLOSE_PAREN))
+	{
+	  gcc_assert (allocators == NULL);
+	  allocators = cur_list;
+	  cur_list = NULL;
+	  break;
+	}
+      else
+	{
+	  cp_parser_error (parser, "expected %<)%>");
+	  cp_parser_skip_to_closing_parenthesis (parser,
+						 /*recovering=*/true,
+						 /*or_comma=*/false,
+						 /*consume_paren=*/true);
+	  goto end;
+	}
+    }
+
+  if (modifiers)
+    for (unsigned i = 0; i < modifiers->length (); i++)
+      {
+	item& it = (*modifiers)[i];
+	const char *p = IDENTIFIER_POINTER (it.name.id);
+	int strcmp_traits = 1, strcmp_memspace = 1;
+
+	if ((strcmp_traits = strcmp ("traits", p)) == 0
+	    || (strcmp_memspace = strcmp ("memspace", p)) == 0)
+	  {
+	    if ((strcmp_traits == 0 && traits_var != NULL_TREE)
+		|| (strcmp_memspace == 0 && memspace_expr != NULL_TREE))
+	      {
+		error_at (it.name.loc, "duplicate %qs modifier", p);
+		goto end;
+	      }
+	    t = cp_parser_lookup_name_simple (parser, it.arg.id, it.arg.loc);
+	    if (t == error_mark_node)
+	      {
+		cp_parser_name_lookup_error (parser, it.arg.id, t, NLE_NULL,
+					     it.arg.loc);
+	      }
+	    else if (strcmp_memspace == 0)
+	      memspace_expr = t;
+	    else if (strcmp_traits == 0)
+	      traits_var = t;
+	    else
+	      gcc_unreachable ();
+	  }
+	else
+	  {
+	    error_at (it.name.loc, "unknown modifier %qE", it.name.id);
+	    goto end;
+	  }
+      }
+
+  if (allocators)
+    {
+      if (modifiers)
+	{
+	  if (allocators->length () > 1)
+	    {
+	      error_at ((*allocators)[1].name.loc,
+			"%<uses_allocators%> clause only accepts a single "
+			"allocator when using modifiers");
+	      goto end;
+	    }
+	  else if ((*allocators)[0].arg.id)
+	    {
+	      error_at ((*allocators)[0].arg.loc,
+			"legacy %<%E(%E)%> traits syntax not allowed in "
+			"%<uses_allocators%> clause when using modifiers",
+			(*allocators)[0].name.id, (*allocators)[0].arg.id);
+	      goto end;
+	    }
+	}
+
+      for (unsigned i = 0; i < allocators->length (); i++)
+	{
+	  item& it = (*allocators)[i];
+	  t = cp_parser_lookup_name_simple (parser, it.name.id, it.name.loc);
+	  if (t == error_mark_node)
+	    {
+	      cp_parser_name_lookup_error (parser, it.name.id, t, NLE_NULL,
+					   it.name.loc);
+	      goto end;
+	    }
+	  else if (t != error_mark_node)
+	    {
+	      tree t2 = NULL_TREE;
+	      if (it.arg.id)
+		{
+		  t2 = cp_parser_lookup_name_simple (parser, it.arg.id,
+						     it.arg.loc);
+		  if (t2 == error_mark_node)
+		    {
+		      cp_parser_name_lookup_error (parser, it.arg.id, t2,
+						   NLE_NULL, it.arg.loc);
+		      goto end;
+		    }
+		}
+	      else
+		t2 = traits_var;
+
+	      tree c = build_omp_clause (clause_loc,
+					 OMP_CLAUSE_USES_ALLOCATORS);
+	      OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c) = t;
+	      OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c) = memspace_expr;
+	      OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c) = t2;
+	      OMP_CLAUSE_CHAIN (c) = nl;
+	      nl = c;
+	    }
+	}
+    }
+ end:
+  if (cur_list)
+    delete cur_list;
+  if (modifiers)
+    delete modifiers;
+  if (allocators)
+    delete allocators;
+  cp_parser_skip_to_closing_parenthesis (parser,
+					 /*recovering=*/false,
+					 /*or_comma=*/false,
+					 /*consume_paren=*/true);
+  return nl;
+}
+
 /* OpenMP 2.5:
    lastprivate ( variable-list )
 
@@ -41856,6 +42086,10 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask,
 	  clauses = cp_parser_omp_clause_allocate (parser, clauses);
 	  c_name = "allocate";
 	  break;
+	case PRAGMA_OMP_CLAUSE_USES_ALLOCATORS:
+	  clauses = cp_parser_omp_clause_uses_allocators (parser, clauses);
+	  c_name = "uses_allocators";
+	  break;
 	case PRAGMA_OMP_CLAUSE_LINEAR:
 	  {
 	    bool declare_simd = false;
@@ -46427,7 +46661,8 @@ cp_parser_omp_target_update (cp_parser *parser, cp_token *pragma_tok,
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IN_REDUCTION)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_THREAD_LIMIT)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR)\
-	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR))
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR)\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USES_ALLOCATORS))
 
 static bool
 cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index 8090c71809f..47796ab68c1 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -8029,6 +8029,101 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	      break;
 	    }
 	  gcc_unreachable ();
+	case OMP_CLAUSE_USES_ALLOCATORS:
+	  t = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c);
+	  if (TREE_CODE (t) == FIELD_DECL)
+	    {
+	      sorry_at (OMP_CLAUSE_LOCATION (c), "class members not yet "
+			"supported in %<uses_allocators%> clause");
+	      remove = true;
+	      break;
+	    }
+	  t = convert_from_reference (t);
+	  if (TREE_CODE (TREE_TYPE (t)) != ENUMERAL_TYPE
+	      || strcmp (IDENTIFIER_POINTER (TYPE_IDENTIFIER (TREE_TYPE (t))),
+			 "omp_allocator_handle_t") != 0)
+	    {
+	      error_at (OMP_CLAUSE_LOCATION (c),
+			"allocator must be of %<omp_allocator_handle_t%> type");
+	      remove = true;
+	      break;
+	    }
+	  if (TREE_CODE (t) == CONST_DECL)
+	    {
+	      /* Currently for pre-defined allocators in libgomp, we do not
+		 require additional init/fini inside target regions, so discard
+		 such clauses.  */
+	      remove = true;
+
+	      if (strcmp (IDENTIFIER_POINTER (DECL_NAME (t)),
+			  "omp_null_allocator") == 0)
+		{
+		  error_at (OMP_CLAUSE_LOCATION (c),
+			    "%<omp_null_allocator%> cannot be used in "
+			    "%<uses_allocators%> clause");
+		  break;
+		}
+
+	      if (OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c)
+		  || OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c))
+		{
+		  error_at (OMP_CLAUSE_LOCATION (c),
+			    "modifiers cannot be used with pre-defined "
+			    "allocators");
+		  break;
+		}
+	    }
+	  t = OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c);
+	  if (t != NULL_TREE
+	      && (TREE_CODE (t) != CONST_DECL
+		  || TREE_CODE (TREE_TYPE (t)) != ENUMERAL_TYPE
+		  || strcmp (IDENTIFIER_POINTER (TYPE_IDENTIFIER (TREE_TYPE (t))),
+			     "omp_memspace_handle_t") != 0))
+	    {
+	      error_at (OMP_CLAUSE_LOCATION (c), "memspace modifier must be "
+			"constant enum of %<omp_memspace_handle_t%> type");
+	      remove = true;
+	      break;
+	    }
+	  t = OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c);
+	  if (t != NULL_TREE)
+	    {
+	      bool type_err = false;
+
+	      if (TREE_CODE (TREE_TYPE (t)) != ARRAY_TYPE
+		  || DECL_SIZE (t) == NULL_TREE)
+		type_err = true;
+	      else
+		{
+		  tree elem_t = TREE_TYPE (TREE_TYPE (t));
+		  if (TREE_CODE (elem_t) != RECORD_TYPE
+		      || strcmp (IDENTIFIER_POINTER (TYPE_IDENTIFIER (elem_t)),
+				 "omp_alloctrait_t") != 0
+		      || !TYPE_READONLY (elem_t))
+		    type_err = true;
+		}
+	      if (type_err)
+		{
+		  error_at (OMP_CLAUSE_LOCATION (c), "traits array %qE must be of "
+			    "%<const omp_alloctrait_t []%> type", t);
+		  remove = true;
+		}
+	      else
+		{
+		  tree cst_val = decl_constant_value (t);
+		  if (cst_val == t)
+		    {
+		      error_at (OMP_CLAUSE_LOCATION (c), "traits array must be "
+				"of constant values");
+
+		      remove = true;
+		    }
+		}
+	    }
+	  if (remove)
+	    break;
+	  pc = &OMP_CLAUSE_CHAIN (c);
+	  continue;
 	case OMP_CLAUSE_DEPEND:
 	case OMP_CLAUSE_AFFINITY:
 	  t = OMP_CLAUSE_DECL (c);
diff --git a/gcc/fortran/trans-array.cc b/gcc/fortran/trans-array.cc
index bbb81f40aa9..3329f8f3729 100644
--- a/gcc/fortran/trans-array.cc
+++ b/gcc/fortran/trans-array.cc
@@ -6501,10 +6501,7 @@ gfc_conv_array_initializer (tree type, gfc_expr * expr)
 			       &expr->where, flag_max_array_constructor);
 	      return NULL_TREE;
 	    }
-          if (mpz_cmp_si (c->offset, 0) != 0)
-            index = gfc_conv_mpz_to_tree (c->offset, gfc_index_integer_kind);
-          else
-            index = NULL_TREE;
+	  index = gfc_conv_mpz_to_tree (c->offset, gfc_index_integer_kind);
 
 	  if (mpz_cmp_si (c->repeat, 1) > 0)
 	    {
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index 82bbc41b388..9905c882d29 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -2758,7 +2758,12 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 			      gfc_init_se (&se, NULL);
 			      gfc_conv_expr (&se, n->u2.allocator);
 			      gfc_add_block_to_block (block, &se.pre);
-			      allocator_ = gfc_evaluate_now (se.expr, block);
+			      t = se.expr;
+			      if (DECL_P (t) && se.post.head == NULL_TREE)
+				allocator_ = (POINTER_TYPE_P (TREE_TYPE (t))
+					      ? build_fold_indirect_ref (t): t);
+			      else
+				allocator_ = gfc_evaluate_now (t, block);
 			      gfc_add_block_to_block (block, &se.post);
 			    }
 			  OMP_CLAUSE_ALLOCATE_ALLOCATOR (node) = allocator_;
@@ -3938,13 +3943,36 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
 	    }
 	  break;
 	case OMP_LIST_USES_ALLOCATORS:
-	  /* Ignore pre-defined allocators as no special treatment is needed. */
 	  for (; n != NULL; n = n->next)
-	    if (n->sym->attr.flavor == FL_VARIABLE)
-	      break;
-	  if (n != NULL)
-	    sorry_at (input_location, "%<uses_allocators%> clause with traits "
-				      "and memory spaces");
+	    {
+	      if (!n->sym->attr.referenced)
+		continue;
+	      tree node = build_omp_clause (input_location,
+					    OMP_CLAUSE_USES_ALLOCATORS);
+	      tree t;
+	      if (n->sym->attr.flavor == FL_VARIABLE)
+		t = gfc_get_symbol_decl (n->sym);
+	      else
+		{
+		  t = gfc_conv_mpz_to_tree (n->sym->value->value.integer,
+					    n->sym->ts.kind);
+		  t = fold_convert (ptr_type_node, t);
+		}
+	      OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR(node) = t;
+	      if (n->u.memspace_sym)
+		{
+		  n->u.memspace_sym->attr.referenced = true;
+		  OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (node)
+		    = gfc_get_symbol_decl (n->u.memspace_sym);
+		}
+	      if (n->u2.traits_sym)
+		{
+		  n->u2.traits_sym->attr.referenced = true;
+		  OMP_CLAUSE_USES_ALLOCATORS_TRAITS (node)
+		    = gfc_get_symbol_decl (n->u2.traits_sym);
+		}
+	      omp_clauses = gfc_trans_add_clause (node, omp_clauses);
+	    }
 	  break;
 	default:
 	  break;
diff --git a/gcc/fortran/types.def b/gcc/fortran/types.def
index 7a465c89c5f..a9374d68d78 100644
--- a/gcc/fortran/types.def
+++ b/gcc/fortran/types.def
@@ -81,6 +81,7 @@ DEF_FUNCTION_TYPE_0 (BT_FN_UINT, BT_UINT)
 DEF_FUNCTION_TYPE_0 (BT_FN_VOID, BT_VOID)
 
 DEF_FUNCTION_TYPE_1 (BT_FN_VOID_PTR, BT_VOID, BT_PTR)
+DEF_FUNCTION_TYPE_1 (BT_FN_VOID_PTRMODE, BT_VOID, BT_PTRMODE)
 DEF_FUNCTION_TYPE_1 (BT_FN_VOID_PTRPTR, BT_VOID, BT_PTR_PTR)
 DEF_FUNCTION_TYPE_1 (BT_FN_VOID_VPTR, BT_VOID, BT_VOLATILE_PTR)
 DEF_FUNCTION_TYPE_1 (BT_FN_INT_INT, BT_INT, BT_INT)
@@ -154,6 +155,8 @@ DEF_FUNCTION_TYPE_3 (BT_FN_VOID_SIZE_SIZE_PTR, BT_VOID, BT_SIZE, BT_SIZE,
 DEF_FUNCTION_TYPE_3 (BT_FN_UINT_UINT_PTR_PTR, BT_UINT, BT_UINT, BT_PTR, BT_PTR)
 DEF_FUNCTION_TYPE_3 (BT_FN_PTR_SIZE_SIZE_PTRMODE,
 		     BT_PTR, BT_SIZE, BT_SIZE, BT_PTRMODE)
+DEF_FUNCTION_TYPE_3 (BT_FN_PTRMODE_PTRMODE_INT_PTR, BT_PTRMODE, BT_PTRMODE,
+		     BT_INT, BT_PTR)
 
 DEF_FUNCTION_TYPE_4 (BT_FN_VOID_OMPFN_PTR_UINT_UINT,
                      BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR, BT_UINT, BT_UINT)
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index d52d71b9b6b..08bd37e2e22 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -1381,18 +1381,46 @@ gimplify_bind_expr (tree *expr_p, gimple_seq *pre_p)
 		 dynamic_allocators clause is present in the same compilation
 		 unit.  */
 	      bool missing_dyn_alloc = false;
-	      if (alloc == NULL_TREE
-		  && ((omp_requires_mask & OMP_REQUIRES_DYNAMIC_ALLOCATORS)
-		      == 0))
+	      if ((omp_requires_mask & OMP_REQUIRES_DYNAMIC_ALLOCATORS) == 0)
 		{
 		  /* This comes too early for omp_discover_declare_target...,
 		     but should at least catch the most common cases.  */
 		  missing_dyn_alloc
-		    = cgraph_node::get (current_function_decl)->offloadable;
+		    = (alloc == NULL_TREE
+		       && cgraph_node::get (current_function_decl)->offloadable);
 		  for (struct gimplify_omp_ctx *ctx2 = ctx;
 		       ctx2 && !missing_dyn_alloc; ctx2 = ctx2->outer_context)
 		    if (ctx2->code == OMP_TARGET)
-		      missing_dyn_alloc = true;
+		      {
+			if (alloc == NULL_TREE)
+			  missing_dyn_alloc = true;
+			else if (TREE_CODE (alloc) != INTEGER_CST)
+			  {
+			    tree alloc2 = alloc;
+			    if (TREE_CODE (alloc2) == MEM_REF
+				|| TREE_CODE (alloc2) == INDIRECT_REF)
+			      alloc2 = TREE_OPERAND (alloc2, 0);
+			    tree c2;
+			    for (c2 = ctx2->clauses; c2;
+				 c2 = OMP_CLAUSE_CHAIN (c2))
+			      if (OMP_CLAUSE_CODE (c2)
+				  == OMP_CLAUSE_USES_ALLOCATORS)
+				{
+				  tree t2
+				    = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c2);
+				  if (operand_equal_p (alloc2, t2))
+				    break;
+				}
+			    if (c2 == NULL_TREE)
+			      error_at (EXPR_LOC_OR_LOC (
+					  alloc, DECL_SOURCE_LOCATION (t)),
+					"%qE in %<allocator%> clause inside a "
+					"target region must be specified in an "
+					"%<uses_allocators%> clause on the "
+					"%<target%> directive", alloc2);
+			  }
+			break;
+		      }
 		}
 	      if (missing_dyn_alloc)
 		error_at (DECL_SOURCE_LOCATION (t),
@@ -12240,6 +12268,21 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	  nowait = 1;
 	  break;
 
+	case OMP_CLAUSE_USES_ALLOCATORS:
+	  if (TREE_CODE (OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c))
+	      != INTEGER_CST)
+	    {
+	      decl = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c);
+	      omp_add_variable (ctx, decl, GOVD_SEEN | GOVD_PRIVATE);
+
+	      decl = OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c);
+	      if (decl && !DECL_INITIAL (decl))
+		omp_add_variable (ctx, decl, GOVD_SEEN | GOVD_FIRSTPRIVATE);
+	    }
+	  else
+	    remove = true;
+	  break;
+
 	case OMP_CLAUSE_ORDERED:
 	case OMP_CLAUSE_UNTIED:
 	case OMP_CLAUSE_COLLAPSE:
@@ -12375,6 +12418,49 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	      remove = true;
 	      break;
 	    }
+	  if ((omp_requires_mask & OMP_REQUIRES_DYNAMIC_ALLOCATORS) == 0
+	      && OMP_CLAUSE_ALLOCATE_ALLOCATOR (c)
+	      && TREE_CODE (OMP_CLAUSE_ALLOCATE_ALLOCATOR (c)) != INTEGER_CST)
+	    {
+	      tree allocator = OMP_CLAUSE_ALLOCATE_ALLOCATOR (c);
+	      tree clauses = NULL_TREE;
+
+	      /* Get clause list of the nearest enclosing target construct.  */
+	      if (ctx->code == OMP_TARGET)
+		clauses = *orig_list_p;
+	      else
+		{
+		  struct gimplify_omp_ctx *tctx = ctx->outer_context;
+		  while (tctx && tctx->code != OMP_TARGET)
+		    tctx = tctx->outer_context;
+		  if (tctx)
+		    clauses = tctx->clauses;
+		}
+
+	      if (clauses)
+		{
+		  tree uc;
+		  if (TREE_CODE (allocator) == MEM_REF
+		      || TREE_CODE (allocator) == INDIRECT_REF)
+		    allocator = TREE_OPERAND (allocator, 0);
+		  for (uc = clauses; uc; uc = OMP_CLAUSE_CHAIN (uc))
+		    if (OMP_CLAUSE_CODE (uc) == OMP_CLAUSE_USES_ALLOCATORS)
+		      {
+			tree uc_allocator
+			  = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (uc);
+			if (operand_equal_p (allocator, uc_allocator))
+			  break;
+		      }
+		  if (uc == NULL_TREE)
+		    {
+		      error_at (OMP_CLAUSE_LOCATION (c), "allocator %qE "
+				"requires %<uses_allocators(%E)%> clause in "
+				"target region", allocator, allocator);
+		      remove = true;
+		      break;
+		    }
+		}
+	    }
 	  if (gimplify_expr (&OMP_CLAUSE_ALLOCATE_ALLOCATOR (c), pre_p, NULL,
 			     is_gimple_val, fb_rvalue) == GS_ERROR)
 	    {
@@ -13411,6 +13497,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 	case OMP_CLAUSE_FINALIZE:
 	case OMP_CLAUSE_INCLUSIVE:
 	case OMP_CLAUSE_EXCLUSIVE:
+	case OMP_CLAUSE_USES_ALLOCATORS:
 	  break;
 
 	case OMP_CLAUSE_NOHOST:
@@ -15777,6 +15864,92 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
 	  body = NULL;
 	  gimple_seq_add_stmt (&body, g);
 	}
+      else if ((ort & ORT_TARGET) != 0 && (ort & ORT_ACC) == 0)
+	{
+	  gimple_seq init_seq = NULL;
+	  gimple_seq fini_seq = NULL;
+
+	  tree omp_init_allocator_fn = NULL_TREE;
+	  tree omp_destroy_allocator_fn = NULL_TREE;
+
+	  for (tree *cp = &OMP_CLAUSES (expr); *cp != NULL;
+	       cp = &OMP_CLAUSE_CHAIN (*cp))
+	    if (OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_USES_ALLOCATORS)
+	      {
+		tree c = *cp;
+		tree allocator = OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c);
+		tree memspace = OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (c);
+		tree traits = OMP_CLAUSE_USES_ALLOCATORS_TRAITS (c);
+
+		if (omp_init_allocator_fn == NULL_TREE)
+		  {
+		    omp_init_allocator_fn
+		      = builtin_decl_explicit (BUILT_IN_OMP_INIT_ALLOCATOR);
+		    omp_destroy_allocator_fn
+		      = builtin_decl_explicit (BUILT_IN_OMP_DESTROY_ALLOCATOR);
+		  }
+		tree ntraits, traits_var;
+		if (traits == NULL_TREE)
+		  {
+		     ntraits = integer_zero_node;
+		     traits_var = null_pointer_node;
+		  }
+		else if (DECL_INITIAL (traits))
+		  {
+		    location_t loc = OMP_CLAUSE_LOCATION (c);
+		    tree t = DECL_INITIAL (traits);
+		    gcc_assert (TREE_CODE (t) == CONSTRUCTOR);
+		    ntraits = build_int_cst (integer_type_node,
+					     CONSTRUCTOR_NELTS (t));
+		    t = get_initialized_tmp_var (t, &init_seq, NULL);
+		    traits_var = build_fold_addr_expr_loc (loc, t);
+		  }
+		else
+		  {
+		    location_t loc = OMP_CLAUSE_LOCATION (c);
+		    gcc_assert (TREE_CODE (TREE_TYPE (traits)) == ARRAY_TYPE);
+		    tree t = TYPE_DOMAIN (TREE_TYPE (traits));
+		    tree min = TYPE_MIN_VALUE (t);
+		    tree max = TYPE_MAX_VALUE (t);
+		    gcc_assert (TREE_CODE (min) == INTEGER_CST
+				&& TREE_CODE (max) == INTEGER_CST);
+		    t = fold_build2_loc (loc, MINUS_EXPR, TREE_TYPE (min),
+					 max, min);
+		    t = fold_build2_loc (loc, PLUS_EXPR, TREE_TYPE (min),
+					 t, build_int_cst (TREE_TYPE (min), 1));
+		    ntraits = t;
+		    traits_var = build_fold_addr_expr_loc (loc, traits);
+		  }
+
+		if (memspace == NULL_TREE)
+		  memspace = build_int_cst (pointer_sized_int_node, 0);
+		else
+		  memspace = fold_convert (pointer_sized_int_node,
+					   memspace);
+
+		tree call = build_call_expr_loc (OMP_CLAUSE_LOCATION (c),
+						 omp_init_allocator_fn, 3,
+						 memspace, ntraits,
+						 traits_var);
+		call = fold_convert (TREE_TYPE (allocator), call);
+		gimplify_assign (allocator, call, &init_seq);
+
+		call = build_call_expr_loc (OMP_CLAUSE_LOCATION (c),
+					    omp_destroy_allocator_fn, 1,
+					    allocator);
+		gimplify_and_add (call, &fini_seq);
+	      }
+
+	  if (fini_seq)
+	    {
+	      gbind *bind = as_a<gbind *> (gimple_seq_first_stmt (body));
+	      g = gimple_build_try (gimple_bind_body (bind),
+				    fini_seq, GIMPLE_TRY_FINALLY);
+	      gimple_seq_add_stmt (&init_seq, g);
+	      gimple_bind_set_body (bind, init_seq);
+	      body = bind;
+	    }
+	}
     }
   else
     gimplify_and_add (OMP_BODY (expr), &body);
diff --git a/gcc/omp-builtins.def b/gcc/omp-builtins.def
index ed78d49d205..c07f8533c0a 100644
--- a/gcc/omp-builtins.def
+++ b/gcc/omp-builtins.def
@@ -76,6 +76,10 @@ DEF_GOMP_BUILTIN (BUILT_IN_OMP_GET_TEAM_NUM, "omp_get_team_num",
 		  BT_FN_INT, ATTR_CONST_NOTHROW_LEAF_LIST)
 DEF_GOMP_BUILTIN (BUILT_IN_OMP_GET_NUM_TEAMS, "omp_get_num_teams",
 		  BT_FN_INT, ATTR_CONST_NOTHROW_LEAF_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_OMP_INIT_ALLOCATOR, "omp_init_allocator",
+		  BT_FN_PTRMODE_PTRMODE_INT_PTR, ATTR_NOTHROW_LEAF_LIST)
+DEF_GOMP_BUILTIN (BUILT_IN_OMP_DESTROY_ALLOCATOR, "omp_destroy_allocator",
+		  BT_FN_VOID_PTRMODE, ATTR_NOTHROW_LEAF_LIST)
 
 DEF_GOMP_BUILTIN (BUILT_IN_GOMP_ATOMIC_START, "GOMP_atomic_start",
 		  BT_FN_VOID, ATTR_NOTHROW_LEAF_LIST)
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 161bcfeec05..a53a620fc1f 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -1166,6 +1166,36 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	    && omp_maybe_offloaded_ctx (ctx))
 	  error_at (OMP_CLAUSE_LOCATION (c), "%<allocate%> clause must"
 		    " specify an allocator here");
+	if ((omp_requires_mask & OMP_REQUIRES_DYNAMIC_ALLOCATORS) == 0
+	    && OMP_CLAUSE_ALLOCATE_ALLOCATOR (c) != NULL_TREE
+	    && DECL_P (OMP_CLAUSE_ALLOCATE_ALLOCATOR (c))
+	    && !DECL_ARTIFICIAL (OMP_CLAUSE_ALLOCATE_ALLOCATOR (c)))
+	  {
+	    tree alloc2 = OMP_CLAUSE_ALLOCATE_ALLOCATOR (c);
+	    if (TREE_CODE (alloc2) == MEM_REF
+		|| TREE_CODE (alloc2) == INDIRECT_REF)
+	      alloc2 = TREE_OPERAND (alloc2, 0);
+	    omp_context *ctx2 = ctx;
+	    for (; ctx2; ctx2 = ctx2->outer)
+	      if (is_gimple_omp_offloaded (ctx2->stmt))
+		break;
+	    if (ctx2 != NULL)
+	      {
+		tree c2 = gimple_omp_target_clauses (ctx2->stmt);
+		for (; c2; c2 = OMP_CLAUSE_CHAIN (c2))
+		  if (OMP_CLAUSE_CODE (c2) == OMP_CLAUSE_USES_ALLOCATORS
+		      && operand_equal_p (
+			   alloc2, OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (c2)))
+		    break;
+		if (c2 == NULL_TREE)
+		  error_at (EXPR_LOC_OR_LOC (OMP_CLAUSE_ALLOCATE_ALLOCATOR (c),
+					     OMP_CLAUSE_LOCATION (c)),
+			    "allocator %qE in %<allocate%> clause inside a "
+			    "target region must be specified in an "
+			    "%<uses_allocators%> clause on the %<target%> "
+			    "directive", alloc2);
+	      }
+	  }
 	if (ctx->allocate_map == NULL)
 	  ctx->allocate_map = new hash_map<tree, tree>;
 	tree val = integer_zero_node;
@@ -1756,6 +1786,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	case OMP_CLAUSE_FINALIZE:
 	case OMP_CLAUSE_TASK_REDUCTION:
 	case OMP_CLAUSE_ALLOCATE:
+	case OMP_CLAUSE_USES_ALLOCATORS:
 	  break;
 
 	case OMP_CLAUSE_ALIGNED:
@@ -1974,6 +2005,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	case OMP_CLAUSE_FINALIZE:
 	case OMP_CLAUSE_FILTER:
 	case OMP_CLAUSE__CONDTEMP_:
+	case OMP_CLAUSE_USES_ALLOCATORS:
 	  break;
 
 	case OMP_CLAUSE__CACHE_:
diff --git a/gcc/testsuite/c-c++-common/gomp/uses_allocators-1.c b/gcc/testsuite/c-c++-common/gomp/uses_allocators-1.c
new file mode 100644
index 00000000000..5a2e4a90e54
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/uses_allocators-1.c
@@ -0,0 +1,46 @@
+typedef enum omp_allocator_handle_t
+#if __cplusplus >= 201103L
+: __UINTPTR_TYPE__
+#endif
+{
+  omp_default_mem_alloc = 1,
+  omp_low_lat_mem_alloc = 5,
+  __omp_allocator_handle_t_max__ = __UINTPTR_MAX__
+} omp_allocator_handle_t;
+
+typedef struct omp_alloctrait_t
+{
+  int key;
+  int value;
+} omp_alloctrait_t;
+
+extern void *omp_alloc (__SIZE_TYPE__, omp_allocator_handle_t);
+
+void
+f (omp_allocator_handle_t my_alloc)
+{
+  #pragma omp target
+  {
+    int a; /* { dg-error "'my_alloc' in 'allocator' clause inside a target region must be specified in an 'uses_allocators' clause on the 'target' directive" "" { target c } } */
+    #pragma omp allocate(a) allocator(my_alloc) /* { dg-message "sorry, unimplemented: '#pragma omp allocate' not yet supported" "" { target c++ } }  */
+    a  = 5;
+    void *prt = omp_alloc(32, my_alloc);
+    #pragma omp parallel allocate(allocator(my_alloc) : a) firstprivate(a) /* { dg-error "allocator 'my_alloc' in 'allocate' clause inside a target region must be specified in an 'uses_allocators' clause on the 'target' directive" } */
+      a = 7;
+  }
+}
+
+void
+g (omp_allocator_handle_t my_alloc)
+{
+  /* The following defines a default-mem-space allocator with no extra traits. */
+  #pragma omp target uses_allocators(my_alloc)
+  {
+    int a;
+    #pragma omp allocate(a) allocator(my_alloc)  /* { dg-message "sorry, unimplemented: '#pragma omp allocate' not yet supported" "" { target c++ } }  */
+    a  = 5;
+    void *prt = omp_alloc(32, my_alloc);
+    #pragma omp parallel allocate(allocator(my_alloc) : a) firstprivate(a)
+      a = 7;
+  }
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/uses_allocators-2.c b/gcc/testsuite/c-c++-common/gomp/uses_allocators-2.c
new file mode 100644
index 00000000000..4dd1f13100a
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/uses_allocators-2.c
@@ -0,0 +1,33 @@
+typedef enum omp_allocator_handle_t
+#if __cplusplus >= 201103L
+: __UINTPTR_TYPE__
+#endif
+{
+  omp_default_mem_alloc = 1,
+  omp_low_lat_mem_alloc = 5,
+  __omp_allocator_handle_t_max__ = __UINTPTR_MAX__
+} omp_allocator_handle_t;
+
+typedef struct omp_alloctrait_t
+{
+  int key;
+  int value;
+} omp_alloctrait_t;
+
+void
+f ()
+{
+   omp_alloctrait_t trait[1] = {{1,1}};
+   omp_allocator_handle_t my_alloc;
+   #pragma omp target uses_allocators(traits(trait) : my_alloc)  /* { dg-error "traits array 'trait' must be of 'const omp_alloctrait_t \\\[\\\]' type" } */
+     ;
+}
+
+void
+g ()
+{
+   const omp_alloctrait_t trait[1] = {{1,1}};
+   omp_allocator_handle_t my_alloc;
+   #pragma omp target uses_allocators(traits(trait) : my_alloc)
+     ;
+}
diff --git a/gcc/testsuite/gfortran.dg/gomp/allocate-1.f90 b/gcc/testsuite/gfortran.dg/gomp/allocate-1.f90
index 8bc6b768778..0463f0e0af9 100644
--- a/gcc/testsuite/gfortran.dg/gomp/allocate-1.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/allocate-1.f90
@@ -24,6 +24,10 @@ module omp_lib_kinds
      parameter :: omp_pteam_mem_alloc = 7
   integer (kind=omp_allocator_handle_kind), &
      parameter :: omp_thread_mem_alloc = 8
+
+  integer, parameter :: omp_memspace_handle_kind = c_intptr_t
+  integer (omp_memspace_handle_kind), &
+     parameter :: omp_default_mem_space = 0
 end module
 
 subroutine bar (a, b, c)
@@ -80,7 +84,8 @@ subroutine foo(x, y)
   
   !$omp target teams distribute parallel do private (x) firstprivate (y) &
   !$omp allocate ((omp_default_mem_alloc + 0):z) allocate &
-  !$omp (omp_default_mem_alloc: x, y) allocate (h: r) lastprivate (z) reduction(+:r)
+  !$omp (omp_default_mem_alloc: x, y) allocate (h: r) lastprivate (z) reduction(+:r) &
+  !$omp uses_allocators(memspace(omp_default_mem_space) : h)
   do i = 1, 10
     call bar (0, x, z);
     call bar2 (1, y, r);
diff --git a/gcc/testsuite/gfortran.dg/gomp/scope-6.f90 b/gcc/testsuite/gfortran.dg/gomp/scope-6.f90
index 4c4f5e034f7..39a65904c33 100644
--- a/gcc/testsuite/gfortran.dg/gomp/scope-6.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/scope-6.f90
@@ -20,4 +20,4 @@ contains
   end
 end
 
-! { dg-final { scan-tree-dump "omp scope private\\(a\\) firstprivate\\(b\\) reduction\\(\\+:c\\) allocate\\(allocator\\(D\\.\[0-9\]+\\):a\\) allocate\\(allocator\\(D\\.\[0-9\]+\\):b\\) allocate\\(allocator\\(D\\.\[0-9\]+\\):c\\)" "original" } }
+! { dg-final { scan-tree-dump "omp scope private\\(a\\) firstprivate\\(b\\) reduction\\(\\+:c\\) allocate\\(allocator\\(h\\):a\\) allocate\\(allocator\\(h\\):b\\) allocate\\(allocator\\(h\\):c\\)" "original" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/uses_allocators-1.f90 b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-1.f90
new file mode 100644
index 00000000000..b20515b3691
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/uses_allocators-1.f90
@@ -0,0 +1,23 @@
+use iso_c_binding
+implicit none
+        integer, parameter :: omp_allocator_handle_kind = c_intptr_t
+        integer, parameter :: omp_alloctrait_key_kind = c_int
+        integer, parameter :: omp_alloctrait_val_kind = c_intptr_t
+        type omp_alloctrait
+          integer (kind=omp_alloctrait_key_kind) key
+          integer (kind=omp_alloctrait_val_kind) value
+        end type omp_alloctrait
+contains
+subroutine x
+integer :: mem
+type(omp_alloctrait), parameter:: mem2(1) = [omp_alloctrait(1,1)]
+integer(omp_allocator_handle_kind) :: var
+!$omp target uses_allocators(memspace(omp_default_mem_space), traits(mem2) : var) defaultmap(none)
+block;
+type(c_ptr) ::c
+c = omp_alloc(omp_default_mem_space, 20_8)
+end block
+!$omp target uses_allocators(omp_default_mem_alloc, var(mem2))
+block; end block
+end
+end
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index 65e51b939a2..c76e7d6116f 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -533,6 +533,9 @@ enum omp_clause_code {
 
   /* OpenACC clause: nohost.  */
   OMP_CLAUSE_NOHOST,
+
+  /* OpenMP clause: uses_allocators.  */
+  OMP_CLAUSE_USES_ALLOCATORS,
 };
 
 #undef DEFTREESTRUCT
diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc
index 1fadd752d05..8cf479d8508 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -800,6 +800,20 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
       pp_right_paren (pp);
       break;
 
+    case OMP_CLAUSE_USES_ALLOCATORS:
+      pp_string (pp, "uses_allocators(");
+      dump_generic_node (pp, OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR (clause),
+			 spc, flags, false);
+      pp_string (pp, ": memspace(");
+      dump_generic_node (pp, OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (clause),
+			 spc, flags, false);
+      pp_string (pp, "), traits(");
+      dump_generic_node (pp, OMP_CLAUSE_USES_ALLOCATORS_TRAITS (clause),
+			 spc, flags, false);
+      pp_right_paren (pp);
+      pp_right_paren (pp);
+      break;
+
     case OMP_CLAUSE_AFFINITY:
       pp_string (pp, "affinity(");
       {
diff --git a/gcc/tree.cc b/gcc/tree.cc
index a3d907a4ac0..46f41cb7167 100644
--- a/gcc/tree.cc
+++ b/gcc/tree.cc
@@ -328,6 +328,7 @@ unsigned const char omp_clause_num_ops[] =
   0, /* OMP_CLAUSE_IF_PRESENT */
   0, /* OMP_CLAUSE_FINALIZE */
   0, /* OMP_CLAUSE_NOHOST */
+  3, /* OMP_CLAUSE_USES_ALLOCATORS */
 };
 
 const char * const omp_clause_code_name[] =
@@ -421,6 +422,7 @@ const char * const omp_clause_code_name[] =
   "if_present",
   "finalize",
   "nohost",
+  "uses_allocators",
 };
 
 /* Unless specific to OpenACC, we tend to internally maintain OpenMP-centric
diff --git a/gcc/tree.h b/gcc/tree.h
index 086b55f0375..88befefe110 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1958,6 +1958,15 @@ class auto_suppress_location_wrappers
 #define OMP_CLAUSE_ALLOCATE_COMBINED(NODE) \
   (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_ALLOCATE)->base.public_flag)
 
+#define OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR(NODE) \
+  OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_USES_ALLOCATORS), 0)
+
+#define OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE(NODE) \
+  OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_USES_ALLOCATORS), 1)
+
+#define OMP_CLAUSE_USES_ALLOCATORS_TRAITS(NODE) \
+  OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_USES_ALLOCATORS), 2)
+
 #define OMP_CLAUSE_NUM_TEAMS_UPPER_EXPR(NODE) \
   OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_NUM_TEAMS), 0)
 
diff --git a/libgomp/testsuite/libgomp.c++/c++.exp b/libgomp/testsuite/libgomp.c++/c++.exp
index ed096e17b9c..5be949bb611 100644
--- a/libgomp/testsuite/libgomp.c++/c++.exp
+++ b/libgomp/testsuite/libgomp.c++/c++.exp
@@ -1,6 +1,15 @@
 load_lib libgomp-dg.exp
 load_gcc_lib gcc-dg.exp
 
+proc check_effective_target_c { } {
+    return 0
+}
+
+proc check_effective_target_c++ { } {
+    return 1
+}
+
+
 if { $blddir != "" } {
     set libstdc++_library_path "../libstdc++-v3/src/.libs"
     set shlib_ext [get_shlib_extension]
diff --git a/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-1.c b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-1.c
new file mode 100644
index 00000000000..eaca686d037
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-1.c
@@ -0,0 +1,53 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" } */
+
+#include <omp.h>
+
+omp_alloctrait_key_t k;
+omp_alloctrait_value_t v;
+
+int main (void)
+{
+  omp_allocator_handle_t foo, bar;
+  const omp_alloctrait_t foo_traits[] = { { omp_atk_pinned,    omp_atv_true },
+  					  { omp_atk_partition, omp_atv_nearest } };
+  #pragma omp target
+    ;
+  #pragma omp target uses_allocators (bar)
+    ;
+  #pragma omp target uses_allocators (foo (foo_traits))
+    ;
+  #pragma omp target uses_allocators (foo (foo_traits), bar (foo_traits))
+    ;
+  #pragma omp target uses_allocators (memspace(omp_high_bw_mem_space) : foo)
+    ;
+  #pragma omp target uses_allocators (traits(foo_traits) : bar)
+    ;
+  #pragma omp target parallel uses_allocators (memspace(omp_high_bw_mem_space), traits(foo_traits) : bar)
+    ;
+  #pragma omp target parallel uses_allocators (traits(foo_traits), memspace(omp_high_bw_mem_space) : bar) uses_allocators(foo)
+  {
+    void *p = omp_alloc ((unsigned long) 32, bar);
+    omp_free (p, bar);
+  }
+  return 0;
+}
+
+/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(bar: memspace\\(\\), traits\\(\\)\\)" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(foo: memspace\\(\\), traits\\(foo_traits\\)\\)" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(bar: memspace\\(\\), traits\\(foo_traits\\)\\) uses_allocators\\(foo: memspace\\(\\), traits\\(foo_traits\\)\\)" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(foo: memspace\\(omp_high_bw_mem_space\\), traits\\(\\)\\)" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(bar: memspace\\(\\), traits\\(foo_traits\\)\\)" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(bar: memspace\\(omp_high_bw_mem_space\\), traits\\(foo_traits\\)\\)" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(bar: memspace\\(omp_high_bw_mem_space\\), traits\\(foo_traits\\)\\) uses_allocators\\(foo: memspace\\(\\), traits\\(\\)\\)" "original" } } */
+
+/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(bar: memspace\\(\\), traits\\(\\)\\) private\\(bar\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(foo: memspace\\(\\), traits\\(foo_traits\\)\\) private\\(foo\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(bar: memspace\\(\\), traits\\(foo_traits\\)\\) uses_allocators\\(foo: memspace\\(\\), traits\\(foo_traits\\)\\) private\\(bar\\) private\\(foo\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(foo: memspace\\(omp_high_bw_mem_space\\), traits\\(\\)\\) private\\(foo\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(bar: memspace\\(\\), traits\\(foo_traits\\)\\) private\\(bar\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(bar: memspace\\(omp_high_bw_mem_space\\), traits\\(foo_traits\\)\\) private\\(bar\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(bar: memspace\\(omp_high_bw_mem_space\\), traits\\(foo_traits\\)\\) uses_allocators\\(foo: memspace\\(\\), traits\\(\\)\\) private\\(bar\\) private\\(foo\\)" "gimple" } } */
+
+/* { dg-final { scan-tree-dump-times "__builtin_omp_init_allocator" 9 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "__builtin_omp_destroy_allocator" 9 "gimple" } } */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-2.c b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-2.c
new file mode 100644
index 00000000000..f350c0a409e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-2.c
@@ -0,0 +1,39 @@
+/* { dg-do compile } */
+
+#include <omp.h>
+
+omp_alloctrait_key_t k;
+omp_alloctrait_value_t v;
+
+int main (void)
+{
+  omp_allocator_handle_t foo, bar;
+  const omp_alloctrait_t traits_array[] = { { omp_atk_pinned,    omp_atv_true },
+					    { omp_atk_partition, omp_atv_nearest } };
+
+  #pragma omp target uses_allocators (baz) /* { dg-error "'baz' undeclared .first use in this function." "" { target c } } */
+    ;                                      /* { dg-error "'baz' has not been declared" "" { target c++ } .-1 } */
+  #pragma omp target uses_allocators (foo (xyz)) /* { dg-error "'xyz' undeclared .first use in this function." "" { target c } } */
+    ;                                            /* { dg-error "'xyz' has not been declared" "" { target c++ } .-1 } */
+  #pragma omp target uses_allocators (foo (traits_array), baz (traits_array)) /* { dg-error "'baz' has not been declared" "" { target c++ } } */
+    ;
+  #pragma omp target uses_allocators (memspace(omp_no_such_space) : foo) /* { dg-error "'omp_no_such_space' undeclared .first use in this function." "" { target c } } */
+    ;                                                                    /* { dg-error "'omp_no_such_space' has not been declared" "" { target c++ } .-1 } */
+  #pragma omp target uses_allocators (memspace(1) : foo) /* { dg-error "expected identifier before numeric constant" } */
+    ;                                                    /* { dg-error "expected '\\\)' before ':' token" "" { target c } .-1 } */
+  #pragma omp target uses_allocators (memspace(omp_no_such_space) : foo, bar) /* { dg-error "'uses_allocators' clause only accepts a single allocator when using modifiers" } */
+    ;                                                                         /* { dg-error "'omp_no_such_space' has not been declared" "" { target c++ } .-1 } */
+  #pragma omp target uses_allocators (traits(xyz) : bar) /* { dg-error "traits array must be of 'const omp_alloctrait_t \\\[\\\]' type" "" { target c } } */
+    ;                                                    /* { dg-error "'xyz' has not been declared" "" { target c++ } .-1 } */
+  #pragma omp target uses_allocators (memspace(omp_high_bw_mem_space), traits(traits_array), memspace (omp_no_such_space) : bar) /* { dg-error "duplicate 'memspace' modifier" } */
+    ;
+  #pragma omp target uses_allocators (traitz(traits_array), memspace(omp_high_bw_mem_space) : bar) /* { dg-error "unknown modifier 'traitz'" } */
+    ;
+  #pragma omp target uses_allocators (omp_null_allocator) /* { dg-error "'omp_null_allocator' cannot be used in 'uses_allocators' clause" } */
+    ;
+  #pragma omp target uses_allocators (memspace(omp_high_bw_mem_space) : foo, bar) /* { dg-error "'uses_allocators' clause only accepts a single allocator when using modifiers" } */
+    ;
+  #pragma omp target uses_allocators (memspace(omp_high_bw_mem_space) : foo(foo_traits)) /* { dg-error "legacy 'foo\\\(foo_traits\\\)' traits syntax not allowed in 'uses_allocators' clause when using modifiers" } */
+    ;
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-3.c b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-3.c
new file mode 100644
index 00000000000..de9ab9213aa
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-3.c
@@ -0,0 +1,37 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" } */
+
+#include <omp.h>
+
+int main (void)
+{
+  omp_allocator_handle_t memspace, traits;
+  const omp_alloctrait_t mytraits[] = { { omp_atk_pinned,    omp_atv_true },
+					{ omp_atk_partition, omp_atv_nearest } };
+  #pragma omp target uses_allocators (memspace)
+    ;
+  #pragma omp target uses_allocators (traits)
+    ;
+  #pragma omp target uses_allocators (traits, memspace)
+    ;
+  #pragma omp target uses_allocators (traits (mytraits))
+    ;
+  #pragma omp target uses_allocators (memspace (mytraits), omp_default_mem_alloc)
+    ;
+  return 0;
+}
+
+/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(memspace: memspace\\(\\), traits\\(\\)\\)" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(traits: memspace\\(\\), traits\\(\\)\\)" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(memspace: memspace\\(\\), traits\\(\\)\\) uses_allocators\\(traits: memspace\\(\\), traits\\(\\)\\)" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(traits: memspace\\(\\), traits\\(mytraits\\)\\)" "original" } } */
+/* { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(memspace: memspace\\(\\), traits\\(mytraits\\)\\)" "original" } } */
+
+/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(memspace: memspace\\(\\), traits\\(\\)\\) private\\(memspace\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(traits: memspace\\(\\), traits\\(\\)\\) private\\(traits\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(memspace: memspace\\(\\), traits\\(\\)\\) uses_allocators\\(traits: memspace\\(\\), traits\\(\\)\\) private\\(traits\\) private\\(memspace\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(traits: memspace\\(\\), traits\\(mytraits\\)\\) private\\(traits\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(memspace: memspace\\(\\), traits\\(mytraits\\)\\) private\\(memspace\\)" "gimple" } } */
+
+/* { dg-final { scan-tree-dump-times "__builtin_omp_init_allocator" 6 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "__builtin_omp_destroy_allocator" 6 "gimple" } } */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-4.c b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-4.c
new file mode 100644
index 00000000000..fabf2a8a319
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/uses_allocators-4.c
@@ -0,0 +1,53 @@
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+#include <stdint.h>
+#include <omp.h>
+
+int
+main ()
+{
+  int x, *xbuf[10];
+  omp_allocator_handle_t my_alloc;
+  const omp_alloctrait_t trait[1]= {{omp_atk_alignment,128}};
+
+  #pragma omp target uses_allocators(omp_low_lat_mem_alloc) map(tofrom: x, xbuf) defaultmap(none)
+    #pragma omp parallel allocate(allocator(omp_low_lat_mem_alloc), align(128): x, xbuf) if(0) firstprivate(x, xbuf)
+      {
+	if ((uintptr_t) &x % 128 != 0)
+	  __builtin_abort ();
+	if ((uintptr_t) xbuf % 128 != 0)
+	  __builtin_abort ();
+      }
+
+  my_alloc = (omp_allocator_handle_t) 0xABCD;
+
+  #pragma omp target uses_allocators(traits(trait): my_alloc) defaultmap(none) map(tofrom: x, xbuf) 
+    #pragma omp parallel allocate(allocator(my_alloc): x, xbuf) if(0) firstprivate(x, xbuf)
+      {
+	if ((uintptr_t) &x % 128 != 0)
+	  __builtin_abort ();
+	if ((uintptr_t) xbuf % 128 != 0)
+	  __builtin_abort ();
+      }
+
+  if (my_alloc != (omp_allocator_handle_t) 0xABCD)
+    __builtin_abort ();
+
+  /* The following creates an allocator with empty traits + default mem space. */
+  #pragma omp target uses_allocators(my_alloc) map(tofrom: x, xbuf) defaultmap(none)
+    #pragma omp parallel allocate(allocator(my_alloc), align(128): x, xbuf) if(0) firstprivate(x, xbuf)
+      {
+	if ((uintptr_t) &x % 128 != 0)
+	  __builtin_abort ();
+	if ((uintptr_t) xbuf % 128 != 0)
+	  __builtin_abort ();
+      }
+
+  if (my_alloc != (omp_allocator_handle_t) 0xABCD)
+    __builtin_abort ();
+
+  return 0;
+}
+
+/* { dg-final { scan-tree-dump-times "#pragma omp target .*private\\(my_alloc\\).*uses_allocators\\(my_alloc: memspace\\(\\), traits\\(trait\\)\\)" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp target .*private\\(my_alloc\\).*uses_allocators\\(my_alloc: memspace\\(\\), traits\\(\\)\\)" 1 "gimple" } } */
diff --git a/libgomp/testsuite/libgomp.c/c.exp b/libgomp/testsuite/libgomp.c/c.exp
index aae282478db..4b59957d1f3 100644
--- a/libgomp/testsuite/libgomp.c/c.exp
+++ b/libgomp/testsuite/libgomp.c/c.exp
@@ -3,6 +3,14 @@ load_gcc_lib gcc-dg.exp
 
 lappend ALWAYS_CFLAGS "compiler=$GCC_UNDER_TEST"
 
+proc check_effective_target_c { } {
+    return 1
+}
+
+proc check_effective_target_c++ { } {
+    return 0
+}
+
 # If a testcase doesn't have special options, use these.
 if ![info exists DEFAULT_CFLAGS] then {
     set DEFAULT_CFLAGS "-O2"
diff --git a/libgomp/testsuite/libgomp.fortran/uses_allocators_2.f90 b/libgomp/testsuite/libgomp.fortran/uses_allocators_2.f90
index 07327969775..bb984033413 100644
--- a/libgomp/testsuite/libgomp.fortran/uses_allocators_2.f90
+++ b/libgomp/testsuite/libgomp.fortran/uses_allocators_2.f90
@@ -3,8 +3,6 @@
 ! Minimal test for valid code:
 ! - predefined allocators do not need any special treatment in uses_allocators
 !   (as 'requires dynamic_allocators' is the default).
-!
-! - Non-predefined allocators are currently rejected ('sorry)'
 
 subroutine test
   use omp_lib
@@ -35,22 +33,22 @@ subroutine non_predef
 
   integer(kind=omp_allocator_handle_kind) :: a1, a2, a3
 
-  !$omp target uses_allocators(omp_default_mem_alloc, a1(trait), a2(trait2))  ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" }
+  !$omp target uses_allocators(omp_default_mem_alloc, a1(trait), a2(trait2))
   block; end block
 
-  !$omp target parallel uses_allocators(omp_default_mem_alloc, a1(trait), a2(trait2))  ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" }
+  !$omp target parallel uses_allocators(omp_default_mem_alloc, a1(trait), a2(trait2))
   block; end block
 
 
   !$omp target uses_allocators(traits(trait):a1) &
-  !$omp&        uses_allocators ( memspace ( omp_low_lat_mem_space ) , traits ( trait2 ) : a2 , a3)  ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" }
+  !$omp&        uses_allocators ( memspace ( omp_low_lat_mem_space ) , traits ( trait2 ) : a2 , a3)
   block; end block
 
   !$omp target parallel uses_allocators(traits(trait):a1) &
-  !$omp&        uses_allocators ( memspace ( omp_low_lat_mem_space ) , traits ( trait2 ) : a2 , a3)  ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" }
+  !$omp&        uses_allocators ( memspace ( omp_low_lat_mem_space ) , traits ( trait2 ) : a2 , a3)
   block; end block
 
-  !$omp target uses_allocators ( traits(trait2) , memspace ( omp_low_lat_mem_space ) : a2 , a3)  ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" }
+  !$omp target uses_allocators ( traits(trait2) , memspace ( omp_low_lat_mem_space ) : a2 , a3)
   block; end block
 end subroutine
 
@@ -62,7 +60,7 @@ subroutine trait_present
   integer(kind=omp_allocator_handle_kind) :: a1
 
   ! Invalid in OpenMP 5.0 / 5.1, but valid since 5.2 the same as omp_default_mem_space + emptry traits array
-  !$omp target uses_allocators ( a1 )  ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" }
+  !$omp target uses_allocators ( a1 )
   block; end block
 end
 
@@ -76,13 +74,13 @@ subroutine odd_names
   integer(kind=omp_allocator_handle_kind) :: traits
   integer(kind=omp_allocator_handle_kind) :: memspace
 
-  !$omp target uses_allocators ( traits(trait1), memspace(trait1) )  ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" }
+  !$omp target uses_allocators ( traits(trait1), memspace(trait1) )
   block; end block
 
-  !$omp target uses_allocators ( traits(trait1), memspace(omp_low_lat_mem_space)  : traits)  ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" }
+  !$omp target uses_allocators ( traits(trait1), memspace(omp_low_lat_mem_space)  : traits)
   block; end block
 
-  !$omp target uses_allocators ( memspace(omp_low_lat_mem_space), traits(trait1) : memspace)  ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" }
+  !$omp target uses_allocators ( memspace(omp_low_lat_mem_space), traits(trait1) : memspace)
   block; end block
 end
 
@@ -94,6 +92,6 @@ subroutine more_checks
   integer(kind=omp_allocator_handle_kind) :: a1, a2(4)
   integer(kind=1) :: a3
 
-  !$omp target uses_allocators(memspace (omp_low_lat_mem_space) : a1 )  ! { dg-message "sorry, unimplemented: 'uses_allocators' clause with traits and memory spaces" }
+  !$omp target uses_allocators(memspace (omp_low_lat_mem_space) : a1 )
   block; end block
 end
diff --git a/libgomp/testsuite/libgomp.fortran/uses_allocators_3.f90 b/libgomp/testsuite/libgomp.fortran/uses_allocators_3.f90
new file mode 100644
index 00000000000..8acdd4223f9
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/uses_allocators_3.f90
@@ -0,0 +1,62 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" }
+
+program main
+  use omp_lib
+  implicit none
+  integer, allocatable :: arr(:)
+  integer (omp_allocator_handle_kind) :: bar, foo
+
+  type (omp_alloctrait), parameter :: traits_array(*) = &
+       [omp_alloctrait(omp_atk_pinned,omp_atv_true),&
+       omp_alloctrait(omp_atk_partition,omp_atv_nearest)]
+
+  !$omp target allocate(bar : arr) uses_allocators(bar)
+  block
+    allocate(arr(100))
+  end block
+
+  !$omp target uses_allocators(omp_default_mem_alloc)
+  block
+  end block
+
+  !$omp target uses_allocators(bar(traits_array), foo (traits_array))
+  block
+    if (foo == 0) stop 1
+  end block
+
+  !$omp target uses_allocators(traits(traits_array) : bar)
+  block
+  end block
+
+  !$omp target parallel uses_allocators(memspace (omp_low_lat_mem_space) : bar)
+  block
+  end block
+
+  !$omp target parallel uses_allocators(memspace (omp_high_bw_mem_space), traits(traits_array) : bar)
+  block
+    use iso_c_binding
+    type(c_ptr) :: ptr
+    integer(c_size_t) :: sz = 32
+    ptr = omp_alloc (sz, bar)
+    call omp_free (ptr, bar)
+  end block
+
+end program main
+
+! { dg-final { scan-tree-dump "pragma omp target allocate\\(allocator\\(bar\\):arr\\) uses_allocators\\(bar: memspace\\(\\), traits\\(\\)\\)" "original" } }
+! { dg-final { scan-tree-dump "pragma omp target" "original" } }
+! { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(bar: memspace\\(\\), traits\\(traits_array\\)\\) uses_allocators\\(foo: memspace\\(\\), traits\\(traits_array\\)\\)" "original" } }
+! { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(bar: memspace\\(\\), traits\\(traits_array\\)\\)" "original" } }
+! { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(bar: memspace\\(omp_low_lat_mem_space\\), traits\\(\\)\\)" "original" } }
+! { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(bar: memspace\\(omp_high_bw_mem_space\\), traits\\(traits_array\\)\\)" "original" } }
+
+! { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) allocate\\(allocator\\(bar\\):arr\\) uses_allocators\\(bar: memspace\\(\\), traits\\(\\)\\) private\\(bar\\)" "gimple" } }
+! { dg-final { scan-tree-dump "pragma omp target" "gimple" } }
+! { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(bar: memspace\\(\\), traits\\(traits_array\\)\\) uses_allocators\\(foo: memspace\\(\\), traits\\(traits_array\\)\\) private\\(foo\\) private\\(bar\\)" "gimple" } }
+! { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(bar: memspace\\(\\), traits\\(traits_array\\)\\) private\\(bar\\)" "gimple" } }
+! { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(bar: memspace\\(omp_low_lat_mem_space\\), traits\\(\\)\\) firstprivate\\(omp_low_lat_mem_space\\) private\\(bar\\)" "gimple" } }
+! { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(bar: memspace\\(omp_high_bw_mem_space\\), traits\\(traits_array\\)\\) firstprivate\\(omp_high_bw_mem_space\\) private\\(bar\\)" "gimple" } }
+
+! { dg-final { scan-tree-dump-times "__builtin_omp_init_allocator" 6 "gimple" } }
+! { dg-final { scan-tree-dump-times "__builtin_omp_destroy_allocator" 6 "gimple" } }
diff --git a/libgomp/testsuite/libgomp.fortran/uses_allocators_4.f90 b/libgomp/testsuite/libgomp.fortran/uses_allocators_4.f90
new file mode 100644
index 00000000000..00f1dcb2763
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/uses_allocators_4.f90
@@ -0,0 +1,54 @@
+! { dg-do compile }
+
+program main
+  use omp_lib
+  implicit none
+  integer (omp_allocator_handle_kind) :: bar, foo
+
+  type (omp_alloctrait), parameter :: traits_array(*) = &
+       [omp_alloctrait(omp_atk_pinned,omp_atv_true),&
+       omp_alloctrait(omp_atk_partition,omp_atv_nearest)]
+
+  !$omp target uses_allocators(omp_non_existant_alloc) ! { dg-error "Allocator 'omp_non_existant_alloc' at .1. in USES_ALLOCATORS must be a scalar integer of kind 'omp_allocator_handle_kind'" }
+  block  ! { dg-error "Symbol 'omp_non_existant_alloc' at .1. has no IMPLICIT type; did you mean 'omp_const_mem_alloc'\?" "" { target *-*-* } .-1 }
+  end block
+
+  !$omp target uses_allocators(bar(traits_array), foo (traits_array), ) ! { dg-error "Invalid character in name" }
+  block
+  end block
+
+  !$omp target uses_allocators(traits(xyz) : bar) ! { dg-error "Symbol 'xyz' at .1. has no IMPLICIT type" }
+  block  ! { dg-error "Traits array 'xyz' in USES_ALLOCATORS .1. must be a one-dimensional named constant array of type 'omp_alloctrait'" "" { target *-*-* } .-1 }
+  end block
+
+  !$omp target uses_allocators(memspace(omp_non_existant_mem_space) : foo) ! { dg-error "Symbol 'omp_non_existant_mem_space' at .1. has no IMPLICIT type; did you mean 'omp_const_mem_space'\?" }
+  ! { dg-error "Memspace 'omp_non_existant_mem_space' at .1. in USES_ALLOCATORS must be a predefined memory space" "" { target *-*-* } .-1 }
+
+  block
+  end block
+
+  !$omp target uses_allocators(traits(traits_array), traits(traits_array) : bar) ! { dg-error "Duplicate TRAITS modifier at .1. in USES_ALLOCATORS clause" }
+  block
+  end block
+
+  !$omp target uses_allocators(memspace(omp_default_mem_space), memspace(omp_default_mem_space) : foo) ! { dg-error "Duplicate MEMSPACE modifier at .1. in USES_ALLOCATORS clause" }
+  block
+  end block
+
+  !$omp target uses_allocators(memspace(omp_default_mem_space), traits(traits_array), traits(traits_array) : foo) ! { dg-error "Duplicate TRAITS modifier at .1. in USES_ALLOCATORS clause" }
+  block
+  end block
+
+  !$omp target uses_allocators (omp_null_allocator) ! { dg-error "Allocator 'omp_null_allocator' at .1. in USES_ALLOCATORS must either a variable or a predefined allocator" }
+  block
+  end block
+
+  !$omp target uses_allocators (memspace(omp_high_bw_mem_space) : foo, bar)
+  block
+  end block
+
+  !$omp target uses_allocators (memspace(omp_high_bw_mem_space) : foo(foo_traits)) ! { dg-error "70:Unexpected '\\(' at .1." }
+  block
+  end block
+
+end program main
diff --git a/libgomp/testsuite/libgomp.fortran/uses_allocators_5.f90 b/libgomp/testsuite/libgomp.fortran/uses_allocators_5.f90
new file mode 100644
index 00000000000..00f87109d2c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/uses_allocators_5.f90
@@ -0,0 +1,14 @@
+! { dg-do compile }
+
+program main
+  use omp_lib
+  implicit none
+  integer, allocatable :: arr(:)
+  integer (omp_allocator_handle_kind) :: bar
+
+  !$omp target allocate(bar : arr) ! { dg-error "allocator 'bar' requires 'uses_allocators.bar.' clause in target region" }
+  block
+    allocate(arr(100))
+  end block
+
+end program main
diff --git a/libgomp/testsuite/libgomp.fortran/uses_allocators_6.f90 b/libgomp/testsuite/libgomp.fortran/uses_allocators_6.f90
new file mode 100644
index 00000000000..3799e3cec73
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/uses_allocators_6.f90
@@ -0,0 +1,50 @@
+! { dg-additional-options "-fdump-tree-gimple" }
+
+program main
+  use iso_c_binding
+  use omp_lib
+  implicit none (type, external)
+  integer :: x, xbuf(10)
+  integer(c_intptr_t) :: iptr
+  integer(omp_allocator_handle_kind) :: my_alloc
+  type(omp_alloctrait), parameter :: trait(*) = [omp_alloctrait(omp_atk_alignment, 128)]
+
+  !$omp target uses_allocators(omp_low_lat_mem_alloc) map(tofrom: x, xbuf) defaultmap(none)
+    !$omp parallel allocate(allocator(omp_low_lat_mem_alloc), align(128): x, xbuf) if(.false.) firstprivate(x, xbuf)
+      if (mod (TRANSFER (loc(x), iptr), 128) /= 0) &
+        stop 1
+      if (mod (TRANSFER (loc(xbuf), iptr), 128) /= 0) &
+        stop 2
+    !$omp end parallel
+  !$omp end target
+
+  my_alloc = transfer(int(z'ABCD', omp_allocator_handle_kind), my_alloc)
+
+  !$omp target uses_allocators(traits(trait): my_alloc) defaultmap(none) map(tofrom: x, xbuf) 
+    !$omp parallel allocate(allocator(my_alloc): x, xbuf) if(.false.) firstprivate(x, xbuf)
+      if (mod (TRANSFER (loc(x), iptr), 128) /= 0) &
+        stop 3
+      if (mod (TRANSFER (loc(xbuf), iptr), 128) /= 0) &
+        stop 4
+    !$omp end parallel
+  !$omp end target
+
+  if (transfer(my_alloc, 0_omp_allocator_handle_kind) /= int(z'ABCD', omp_allocator_handle_kind)) &
+    stop 5
+
+  ! The following creates an allocator with empty traits + default mem space.
+  !$omp target uses_allocators(my_alloc) map(tofrom: x, xbuf) defaultmap(none)
+    !$omp parallel allocate(allocator(my_alloc), align(128): x, xbuf) if(.false.) firstprivate(x, xbuf)
+      if (mod (TRANSFER (loc(x), iptr), 128) /= 0) &
+        stop 6
+      if (mod (TRANSFER (loc(xbuf), iptr), 128) /= 0) &
+        stop 7
+    !$omp end parallel
+  !$omp end target
+
+  if (transfer(my_alloc, 0_omp_allocator_handle_kind) /= int(z'ABCD', omp_allocator_handle_kind)) &
+    stop 8
+end
+
+! { dg-final { scan-tree-dump-times "#pragma omp target .*private\\(my_alloc\\).*uses_allocators\\(my_alloc: memspace\\(\\), traits\\(trait\\)\\)" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target .*private\\(my_alloc\\).*uses_allocators\\(my_alloc: memspace\\(\\), traits\\(\\)\\)" 1 "gimple" } }
