From 8aead58dfcd599634ffb831526046923018d0cd1 Mon Sep 17 00:00:00 2001
From: Alexey Bataev <a.bataev@hotmail.com>
Date: Mon, 1 Apr 2019 16:56:59 +0000
Subject: [PATCH] [OPENMP]Allocate clause allocator in target region.

According to 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. Patch adds a check for this restriction.

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@357412 91177308-0d34-0410-b5e6-96231b3b80d8
---
 include/clang/Basic/DiagnosticSemaKinds.td | 3 +++
 lib/Sema/SemaOpenMP.cpp                    | 9 +++++++++
 test/OpenMP/nvptx_allocate_messages.cpp    | 6 ++++++
 3 files changed, 18 insertions(+)

diff --git a/include/clang/Basic/DiagnosticSemaKinds.td b/include/clang/Basic/DiagnosticSemaKinds.td
index c7818b3092..6b824da360 100644
--- a/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/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/lib/Sema/SemaOpenMP.cpp b/lib/Sema/SemaOpenMP.cpp
index d564bf6bad..9a9bdfaaf5 100644
--- a/lib/Sema/SemaOpenMP.cpp
+++ b/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/test/OpenMP/nvptx_allocate_messages.cpp b/test/OpenMP/nvptx_allocate_messages.cpp
index 5bb9422227..e6fb83f634 100644
--- a/test/OpenMP/nvptx_allocate_messages.cpp
+++ b/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());
-- 
2.40.0