3 files changed, 18 insertions, 0 deletions
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index c7818b30929..6b824da3602 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -9158,6 +9158,9 @@ def note_omp_previous_allocator : Note<
 def err_expected_allocator_clause : Error<"expected an 'allocator' clause "
   "inside of the target region; provide an 'allocator' clause or use 'requires'"
   " directive with the 'dynamic_allocators' clause">;
+def err_expected_allocator_expression : Error<"expected an allocator expression "
+  "inside of the target region; provide an allocator expression or use 'requires'"
+  " directive with the 'dynamic_allocators' clause">;
 def warn_omp_allocate_thread_on_task_target_directive : Warning<
   "allocator with the 'thread' trait access has unspecified behavior on '%0' directive">,
   InGroup<OpenMPClauses>;
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index d564bf6bad4..9a9bdfaaf5d 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -14852,6 +14852,15 @@ OMPClause *Sema::ActOnOpenMPAllocateClause(
     if (AllocatorRes.isInvalid())
       return nullptr;
     Allocator = AllocatorRes.get();
+  } else {
+    // OpenMP 5.0, 2.11.4 allocate Clause, Restrictions.
+    // allocate clauses that appear on a target construct or on constructs in a
+    // target region must specify an allocator expression unless a requires
+    // directive with the dynamic_allocators clause is present in the same
+    // compilation unit.
+    if (LangOpts.OpenMPIsDevice &&
+        !DSAStack->hasRequiresDeclWithClause<OMPDynamicAllocatorsClause>())
+      targetDiag(StartLoc, diag::err_expected_allocator_expression);
   }
   // Analyze and build list of variables.
   SmallVector<Expr *, 8> Vars;
diff --git a/clang/test/OpenMP/nvptx_allocate_messages.cpp b/clang/test/OpenMP/nvptx_allocate_messages.cpp
index 5bb94222270..e6fb83f6346 100644
--- a/clang/test/OpenMP/nvptx_allocate_messages.cpp
+++ b/clang/test/OpenMP/nvptx_allocate_messages.cpp
@@ -57,6 +57,11 @@ template <class T> T foo() {
   T v;
   #pragma omp allocate(v) allocator(omp_cgroup_mem_alloc)
   v = ST<T>::m;
+#if defined(DEVICE) && !defined(REQUIRES)
+// expected-error@+2 2 {{expected an allocator expression inside of the target region; provide an allocator expression or use 'requires' directive with the 'dynamic_allocators' clause}}
+#endif // DEVICE && !REQUIRES
+#pragma omp parallel private(v) allocate(v)
+  v = 0;
   return v;
 }
 
@@ -75,6 +80,7 @@ int main () {
 #endif // DEVICE && !REQUIRES
 #pragma omp allocate(b)
 #if defined(DEVICE) && !defined(REQUIRES)
+// expected-note@+3 {{in instantiation of function template specialization 'foo<int>' requested here}}
 // expected-note@+2 {{called by 'main'}}
 #endif // DEVICE && !REQUIRES
   return (foo<int>() + bar());