diff options
| -rw-r--r-- | clang/include/clang/Basic/DiagnosticSemaKinds.td | 3 | ||||
| -rw-r--r-- | clang/lib/Sema/SemaOpenMP.cpp | 9 | ||||
| -rw-r--r-- | clang/test/OpenMP/nvptx_allocate_messages.cpp | 6 |
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()); |
