r369775 - [OPENMP5.0]Add support for device_type clause in declare target
Alexey Bataev via cfe-commits
cfe-commits at lists.llvm.org
Fri Aug 23 09:11:14 PDT 2019
Author: abataev
Date: Fri Aug 23 09:11:14 2019
New Revision: 369775
URL: http://llvm.org/viewvc/llvm-project?rev=369775&view=rev
Log:
[OPENMP5.0]Add support for device_type clause in declare target
construct.
OpenMP 5.0 introduced new clause for declare target directive, device_type clause, which may accept values host, nohost, and any. Host means
that the function must be emitted only for the host, nohost - only for
the device, and any - for both, device and the host.
Modified:
cfe/trunk/include/clang/AST/RecursiveASTVisitor.h
cfe/trunk/include/clang/Basic/Attr.td
cfe/trunk/include/clang/Basic/AttrDocs.td
cfe/trunk/include/clang/Basic/DiagnosticParseKinds.td
cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
cfe/trunk/include/clang/Basic/OpenMPKinds.def
cfe/trunk/include/clang/Basic/OpenMPKinds.h
cfe/trunk/include/clang/Sema/Sema.h
cfe/trunk/lib/AST/ASTTypeTraits.cpp
cfe/trunk/lib/AST/OpenMPClause.cpp
cfe/trunk/lib/Basic/OpenMPKinds.cpp
cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp
cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp
cfe/trunk/lib/CodeGen/CodeGenModule.cpp
cfe/trunk/lib/Parse/ParseOpenMP.cpp
cfe/trunk/lib/Sema/Sema.cpp
cfe/trunk/lib/Sema/SemaExpr.cpp
cfe/trunk/lib/Sema/SemaOpenMP.cpp
cfe/trunk/lib/Serialization/ASTReaderDecl.cpp
cfe/trunk/test/OpenMP/declare_target_ast_print.cpp
cfe/trunk/test/OpenMP/declare_target_codegen.cpp
cfe/trunk/test/OpenMP/declare_target_messages.cpp
cfe/trunk/test/OpenMP/nvptx_allocate_messages.cpp
cfe/trunk/test/OpenMP/nvptx_asm_delayed_diags.c
cfe/trunk/test/OpenMP/target_vla_messages.cpp
Modified: cfe/trunk/include/clang/AST/RecursiveASTVisitor.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/AST/RecursiveASTVisitor.h?rev=369775&r1=369774&r2=369775&view=diff
==============================================================================
--- cfe/trunk/include/clang/AST/RecursiveASTVisitor.h (original)
+++ cfe/trunk/include/clang/AST/RecursiveASTVisitor.h Fri Aug 23 09:11:14 2019
@@ -2843,6 +2843,7 @@ bool RecursiveASTVisitor<Derived>::Trave
#include "clang/Basic/OpenMPKinds.def"
case OMPC_threadprivate:
case OMPC_uniform:
+ case OMPC_device_type:
case OMPC_unknown:
break;
}
Modified: cfe/trunk/include/clang/Basic/Attr.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/Attr.td?rev=369775&r1=369774&r2=369775&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/Attr.td (original)
+++ cfe/trunk/include/clang/Basic/Attr.td Fri Aug 23 09:11:14 2019
@@ -3207,11 +3207,16 @@ def OMPDeclareTargetDecl : InheritableAt
let Args = [
EnumArgument<"MapType", "MapTypeTy",
[ "to", "link" ],
- [ "MT_To", "MT_Link" ]>
+ [ "MT_To", "MT_Link" ]>,
+ EnumArgument<"DevType", "DevTypeTy",
+ [ "host", "nohost", "any" ],
+ [ "DT_Host", "DT_NoHost", "DT_Any" ]>
];
let AdditionalMembers = [{
void printPrettyPragma(raw_ostream &OS, const PrintingPolicy &Policy) const {
// Use fake syntax because it is for testing and debugging purpose only.
+ if (getDevType() != DT_Any)
+ OS << " device_type(" << ConvertDevTypeTyToStr(getDevType()) << ")";
if (getMapType() != MT_To)
OS << ' ' << ConvertMapTypeTyToStr(getMapType());
}
@@ -3224,6 +3229,14 @@ def OMPDeclareTargetDecl : InheritableAt
return llvm::None;
}
+ static llvm::Optional<DevTypeTy> getDeviceType(const ValueDecl *VD) {
+ if (!VD->hasAttrs())
+ return llvm::None;
+ if (const auto *Attr = VD->getAttr<OMPDeclareTargetDeclAttr>())
+ return Attr->getDevType();
+
+ return llvm::None;
+ }
}];
}
Modified: cfe/trunk/include/clang/Basic/AttrDocs.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/AttrDocs.td?rev=369775&r1=369774&r2=369775&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/AttrDocs.td (original)
+++ cfe/trunk/include/clang/Basic/AttrDocs.td Fri Aug 23 09:11:14 2019
@@ -3176,6 +3176,27 @@ The syntax of the declare target directi
#pragma omp declare target new-line
declarations-definition-seq
#pragma omp end declare target new-line
+
+or
+
+ .. code-block:: c
+
+ #pragma omp declare target (extended-list) new-line
+
+or
+
+ .. code-block:: c
+
+ #pragma omp declare target clause[ [,] clause ... ] new-line
+
+where clause is one of the following:
+
+
+ .. code-block:: c
+
+ to(extended-list)
+ link(list)
+ device_type(host | nohost | any)
}];
}
Modified: cfe/trunk/include/clang/Basic/DiagnosticParseKinds.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticParseKinds.td?rev=369775&r1=369774&r2=369775&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/DiagnosticParseKinds.td (original)
+++ cfe/trunk/include/clang/Basic/DiagnosticParseKinds.td Fri Aug 23 09:11:14 2019
@@ -1195,13 +1195,16 @@ def err_omp_declare_simd_inbranch_notinb
def err_expected_end_declare_target : Error<
"expected '#pragma omp end declare target'">;
def err_omp_declare_target_unexpected_clause: Error<
- "unexpected '%0' clause, only 'to' or 'link' clauses expected">;
+ "unexpected '%0' clause, only %select{'to' or 'link'|'to', 'link' or 'device_type'}1 clauses expected">;
def err_omp_expected_clause: Error<
"expected at least one clause on '#pragma omp %0' directive">;
def err_omp_mapper_illegal_identifier : Error<
"illegal OpenMP user-defined mapper identifier">;
def err_omp_mapper_expected_declarator : Error<
"expected declarator on 'omp declare mapper' directive">;
+def warn_omp_more_one_device_type_clause : Warning<
+ "more than one 'device_type' clause is specified">,
+ InGroup<OpenMPClauses>;
// Pragma loop support.
def err_pragma_loop_missing_argument : Error<
Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td?rev=369775&r1=369774&r2=369775&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td (original)
+++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td Fri Aug 23 09:11:14 2019
@@ -9331,6 +9331,14 @@ def err_omp_unsupported_type : Error <
"host requires %0 bit size %1 type support, but device '%2' does not support it">;
def err_omp_lambda_capture_in_declare_target_not_to : Error<
"variable captured in declare target region must appear in a to clause">;
+def err_omp_device_type_mismatch : Error<
+ "'device_type(%0)' does not match previously specified 'device_type(%1)' for the same declaration">;
+def err_omp_wrong_device_function_call : Error<
+ "function with 'device_type(%0)' is not available on %select{device|host}1">;
+def note_omp_marked_device_type_here : Note<"marked as 'device_type(%0)' here">;
+def warn_omp_declare_target_after_first_use : Warning<
+ "declaration marked as declare target after first use, it may lead to incorrect results">,
+ InGroup<OpenMPTarget>;
} // end of OpenMP category
let CategoryName = "Related Result Type Issue" in {
Modified: cfe/trunk/include/clang/Basic/OpenMPKinds.def
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/OpenMPKinds.def?rev=369775&r1=369774&r2=369775&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/OpenMPKinds.def (original)
+++ cfe/trunk/include/clang/Basic/OpenMPKinds.def Fri Aug 23 09:11:14 2019
@@ -191,6 +191,9 @@
#ifndef OPENMP_ALLOCATE_CLAUSE
# define OPENMP_ALLOCATE_CLAUSE(Name)
#endif
+#ifndef OPENMP_DEVICE_TYPE_KIND
+#define OPENMP_DEVICE_TYPE_KIND(Name)
+#endif
// OpenMP directives.
OPENMP_DIRECTIVE(threadprivate)
@@ -950,6 +953,12 @@ OPENMP_TASKGROUP_CLAUSE(allocate)
// Clauses allowed for OpenMP directive 'declare mapper'.
OPENMP_DECLARE_MAPPER_CLAUSE(map)
+// Device types for 'device_type' clause.
+OPENMP_DEVICE_TYPE_KIND(host)
+OPENMP_DEVICE_TYPE_KIND(nohost)
+OPENMP_DEVICE_TYPE_KIND(any)
+
+#undef OPENMP_DEVICE_TYPE_KIND
#undef OPENMP_ALLOCATE_CLAUSE
#undef OPENMP_DECLARE_MAPPER_CLAUSE
#undef OPENMP_TASKGROUP_CLAUSE
Modified: cfe/trunk/include/clang/Basic/OpenMPKinds.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/OpenMPKinds.h?rev=369775&r1=369774&r2=369775&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/OpenMPKinds.h (original)
+++ cfe/trunk/include/clang/Basic/OpenMPKinds.h Fri Aug 23 09:11:14 2019
@@ -35,6 +35,7 @@ enum OpenMPClauseKind {
#include "clang/Basic/OpenMPKinds.def"
OMPC_threadprivate,
OMPC_uniform,
+ OMPC_device_type,
OMPC_unknown
};
@@ -152,6 +153,14 @@ enum OpenMPAtomicDefaultMemOrderClauseKi
OMPC_ATOMIC_DEFAULT_MEM_ORDER_unknown
};
+/// OpenMP device type for 'device_type' clause.
+enum OpenMPDeviceType {
+#define OPENMP_DEVICE_TYPE_KIND(Name) \
+ OMPC_DEVICE_TYPE_##Name,
+#include "clang/Basic/OpenMPKinds.def"
+ OMPC_DEVICE_TYPE_unknown
+};
+
/// Scheduling data for loop-based OpenMP directives.
struct OpenMPScheduleTy final {
OpenMPScheduleClauseKind Schedule = OMPC_SCHEDULE_unknown;
Modified: cfe/trunk/include/clang/Sema/Sema.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Sema/Sema.h?rev=369775&r1=369774&r2=369775&view=diff
==============================================================================
--- cfe/trunk/include/clang/Sema/Sema.h (original)
+++ cfe/trunk/include/clang/Sema/Sema.h Fri Aug 23 09:11:14 2019
@@ -9001,10 +9001,18 @@ private:
void checkOpenMPDeviceFunction(SourceLocation Loc, FunctionDecl *Callee,
bool CheckForDelayedContext = true);
+ /// Check whether we're allowed to call Callee from the current function.
+ void checkOpenMPHostFunction(SourceLocation Loc, FunctionDecl *Callee,
+ bool CheckCaller = true);
+
/// Check if the expression is allowed to be used in expressions for the
/// OpenMP devices.
void checkOpenMPDeviceExpr(const Expr *E);
+ /// Finishes analysis of the deferred functions calls that may be declared as
+ /// host/nohost during device/host compilation.
+ void finalizeOpenMPDelayedAnalysis();
+
/// Checks if a type or a declaration is disabled due to the owning extension
/// being disabled, and emits diagnostic messages if it is disabled.
/// \param D type or declaration to be checked.
@@ -9151,11 +9159,16 @@ public:
bool ActOnStartOpenMPDeclareTargetDirective(SourceLocation Loc);
/// Called at the end of target region i.e. '#pragme omp end declare target'.
void ActOnFinishOpenMPDeclareTargetDirective();
+ /// Searches for the provided declaration name for OpenMP declare target
+ /// directive.
+ NamedDecl *
+ lookupOpenMPDeclareTargetName(Scope *CurScope, CXXScopeSpec &ScopeSpec,
+ const DeclarationNameInfo &Id,
+ NamedDeclSetType &SameDirectiveDecls);
/// Called on correct id-expression from the '#pragma omp declare target'.
- void ActOnOpenMPDeclareTargetName(Scope *CurScope, CXXScopeSpec &ScopeSpec,
- const DeclarationNameInfo &Id,
+ void ActOnOpenMPDeclareTargetName(NamedDecl *ND, SourceLocation Loc,
OMPDeclareTargetDeclAttr::MapTypeTy MT,
- NamedDeclSetType &SameDirectiveDecls);
+ OMPDeclareTargetDeclAttr::DevTypeTy DT);
/// Check declaration inside target region.
void
checkDeclIsAllowedInOpenMPTarget(Expr *E, Decl *D,
@@ -10539,6 +10552,21 @@ public:
/// // Otherwise, continue parsing as normal.
DeviceDiagBuilder diagIfOpenMPDeviceCode(SourceLocation Loc, unsigned DiagID);
+ /// Creates a DeviceDiagBuilder that emits the diagnostic if the current
+ /// context is "used as host code".
+ ///
+ /// - If CurContext is a `declare target` function or it is known that the
+ /// function is emitted for the host, emits the diagnostics immediately.
+ /// - If CurContext is a non-host function, just ignore it.
+ ///
+ /// Example usage:
+ ///
+ /// // Variable-length arrays are not allowed in NVPTX device code.
+ /// if (diagIfOpenMPHostode(Loc, diag::err_vla_unsupported))
+ /// return ExprError();
+ /// // Otherwise, continue parsing as normal.
+ DeviceDiagBuilder diagIfOpenMPHostCode(SourceLocation Loc, unsigned DiagID);
+
DeviceDiagBuilder targetDiag(SourceLocation Loc, unsigned DiagID);
enum CUDAFunctionTarget {
Modified: cfe/trunk/lib/AST/ASTTypeTraits.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/ASTTypeTraits.cpp?rev=369775&r1=369774&r2=369775&view=diff
==============================================================================
--- cfe/trunk/lib/AST/ASTTypeTraits.cpp (original)
+++ cfe/trunk/lib/AST/ASTTypeTraits.cpp Fri Aug 23 09:11:14 2019
@@ -116,6 +116,7 @@ ASTNodeKind ASTNodeKind::getFromNode(con
#include "clang/Basic/OpenMPKinds.def"
case OMPC_threadprivate:
case OMPC_uniform:
+ case OMPC_device_type:
case OMPC_unknown:
llvm_unreachable("unexpected OpenMP clause kind");
}
Modified: cfe/trunk/lib/AST/OpenMPClause.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/OpenMPClause.cpp?rev=369775&r1=369774&r2=369775&view=diff
==============================================================================
--- cfe/trunk/lib/AST/OpenMPClause.cpp (original)
+++ cfe/trunk/lib/AST/OpenMPClause.cpp Fri Aug 23 09:11:14 2019
@@ -43,6 +43,7 @@ OMPClause::child_range OMPClause::used_c
#include "clang/Basic/OpenMPKinds.def"
case OMPC_threadprivate:
case OMPC_uniform:
+ case OMPC_device_type:
case OMPC_unknown:
break;
}
@@ -127,6 +128,7 @@ const OMPClauseWithPreInit *OMPClauseWit
case OMPC_reverse_offload:
case OMPC_dynamic_allocators:
case OMPC_atomic_default_mem_order:
+ case OMPC_device_type:
break;
}
@@ -203,6 +205,7 @@ const OMPClauseWithPostUpdate *OMPClause
case OMPC_reverse_offload:
case OMPC_dynamic_allocators:
case OMPC_atomic_default_mem_order:
+ case OMPC_device_type:
break;
}
Modified: cfe/trunk/lib/Basic/OpenMPKinds.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/OpenMPKinds.cpp?rev=369775&r1=369774&r2=369775&view=diff
==============================================================================
--- cfe/trunk/lib/Basic/OpenMPKinds.cpp (original)
+++ cfe/trunk/lib/Basic/OpenMPKinds.cpp Fri Aug 23 09:11:14 2019
@@ -55,6 +55,7 @@ OpenMPClauseKind clang::getOpenMPClauseK
#define OPENMP_CLAUSE(Name, Class) .Case(#Name, OMPC_##Name)
#include "clang/Basic/OpenMPKinds.def"
.Case("uniform", OMPC_uniform)
+ .Case("device_type", OMPC_device_type)
.Default(OMPC_unknown);
}
@@ -71,6 +72,8 @@ const char *clang::getOpenMPClauseName(O
return "uniform";
case OMPC_threadprivate:
return "threadprivate or thread local";
+ case OMPC_device_type:
+ return "device_type";
}
llvm_unreachable("Invalid OpenMP clause kind");
}
@@ -145,6 +148,11 @@ unsigned clang::getOpenMPSimpleClauseTyp
.Case(#Name, OMPC_ATOMIC_DEFAULT_MEM_ORDER_##Name)
#include "clang/Basic/OpenMPKinds.def"
.Default(OMPC_ATOMIC_DEFAULT_MEM_ORDER_unknown);
+ case OMPC_device_type:
+ return llvm::StringSwitch<OpenMPDeviceType>(Str)
+#define OPENMP_DEVICE_TYPE_KIND(Name) .Case(#Name, OMPC_DEVICE_TYPE_##Name)
+#include "clang/Basic/OpenMPKinds.def"
+ .Default(OMPC_DEVICE_TYPE_unknown);
case OMPC_unknown:
case OMPC_threadprivate:
case OMPC_if:
@@ -328,6 +336,16 @@ const char *clang::getOpenMPSimpleClause
#include "clang/Basic/OpenMPKinds.def"
}
llvm_unreachable("Invalid OpenMP 'atomic_default_mem_order' clause type");
+ case OMPC_device_type:
+ switch (Type) {
+ case OMPC_DEVICE_TYPE_unknown:
+ return "unknown";
+#define OPENMP_DEVICE_TYPE_KIND(Name) \
+ case OMPC_DEVICE_TYPE_##Name: \
+ return #Name;
+#include "clang/Basic/OpenMPKinds.def"
+ }
+ llvm_unreachable("Invalid OpenMP 'device_type' clause type");
case OMPC_unknown:
case OMPC_threadprivate:
case OMPC_if:
Modified: cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp?rev=369775&r1=369774&r2=369775&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp Fri Aug 23 09:11:14 2019
@@ -9604,14 +9604,28 @@ void CGOpenMPRuntime::scanForTargetRegio
bool CGOpenMPRuntime::emitTargetFunctions(GlobalDecl GD) {
// If emitting code for the host, we do not process FD here. Instead we do
// the normal code generation.
- if (!CGM.getLangOpts().OpenMPIsDevice)
+ if (!CGM.getLangOpts().OpenMPIsDevice) {
+ if (const auto *FD = dyn_cast<FunctionDecl>(GD.getDecl())) {
+ Optional<OMPDeclareTargetDeclAttr::DevTypeTy> DevTy =
+ OMPDeclareTargetDeclAttr::getDeviceType(FD);
+ // Do not emit device_type(nohost) functions for the host.
+ if (DevTy && *DevTy == OMPDeclareTargetDeclAttr::DT_NoHost)
+ return true;
+ }
return false;
+ }
const ValueDecl *VD = cast<ValueDecl>(GD.getDecl());
StringRef Name = CGM.getMangledName(GD);
// Try to detect target regions in the function.
- if (const auto *FD = dyn_cast<FunctionDecl>(VD))
+ if (const auto *FD = dyn_cast<FunctionDecl>(VD)) {
scanForTargetRegionsFunctions(FD->getBody(), Name);
+ Optional<OMPDeclareTargetDeclAttr::DevTypeTy> DevTy =
+ OMPDeclareTargetDeclAttr::getDeviceType(FD);
+ // Do not emit device_type(nohost) functions for the host.
+ if (DevTy && *DevTy == OMPDeclareTargetDeclAttr::DT_Host)
+ return true;
+ }
// Do not to emit function if it is not marked as declare target.
return !OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD) &&
Modified: cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp?rev=369775&r1=369774&r2=369775&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp Fri Aug 23 09:11:14 2019
@@ -4021,6 +4021,7 @@ static void emitOMPAtomicExpr(CodeGenFun
case OMPC_reverse_offload:
case OMPC_dynamic_allocators:
case OMPC_atomic_default_mem_order:
+ case OMPC_device_type:
llvm_unreachable("Clause is not allowed in 'omp atomic'.");
}
}
Modified: cfe/trunk/lib/CodeGen/CodeGenModule.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CodeGenModule.cpp?rev=369775&r1=369774&r2=369775&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CodeGenModule.cpp (original)
+++ cfe/trunk/lib/CodeGen/CodeGenModule.cpp Fri Aug 23 09:11:14 2019
@@ -2122,6 +2122,10 @@ void CodeGenModule::EmitDeferred() {
if (!GV->isDeclaration())
continue;
+ // If this is OpenMP, check if it is legal to emit this global normally.
+ if (LangOpts.OpenMP && OpenMPRuntime && OpenMPRuntime->emitTargetGlobal(D))
+ continue;
+
// Otherwise, emit the definition and move on to the next one.
EmitGlobalDefinition(D, GV);
@@ -2318,11 +2322,20 @@ bool CodeGenModule::MustBeEmitted(const
}
bool CodeGenModule::MayBeEmittedEagerly(const ValueDecl *Global) {
- if (const auto *FD = dyn_cast<FunctionDecl>(Global))
+ if (const auto *FD = dyn_cast<FunctionDecl>(Global)) {
if (FD->getTemplateSpecializationKind() == TSK_ImplicitInstantiation)
// Implicit template instantiations may change linkage if they are later
// explicitly instantiated, so they should not be emitted eagerly.
return false;
+ // In OpenMP 5.0 function may be marked as device_type(nohost) and we should
+ // not emit them eagerly unless we sure that the function must be emitted on
+ // the host.
+ if (LangOpts.OpenMP >= 50 && !LangOpts.OpenMPSimd &&
+ !LangOpts.OpenMPIsDevice &&
+ !OMPDeclareTargetDeclAttr::getDeviceType(FD) &&
+ !FD->isUsed(/*CheckUsedAttr=*/false) && !FD->isReferenced())
+ return false;
+ }
if (const auto *VD = dyn_cast<VarDecl>(Global))
if (Context.getInlineVariableDefinitionKind(VD) ==
ASTContext::InlineVariableDefinitionKind::WeakUnknown)
@@ -2445,8 +2458,7 @@ void CodeGenModule::EmitGlobal(GlobalDec
}
if (LangOpts.OpenMP) {
- // If this is OpenMP device, check if it is legal to emit this global
- // normally.
+ // If this is OpenMP, check if it is legal to emit this global normally.
if (OpenMPRuntime && OpenMPRuntime->emitTargetGlobal(GD))
return;
if (auto *DRD = dyn_cast<OMPDeclareReductionDecl>(Global)) {
Modified: cfe/trunk/lib/Parse/ParseOpenMP.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Parse/ParseOpenMP.cpp?rev=369775&r1=369774&r2=369775&view=diff
==============================================================================
--- cfe/trunk/lib/Parse/ParseOpenMP.cpp (original)
+++ cfe/trunk/lib/Parse/ParseOpenMP.cpp Fri Aug 23 09:11:14 2019
@@ -782,25 +782,114 @@ Parser::ParseOMPDeclareSimdClauses(Parse
LinModifiers, Steps, SourceRange(Loc, EndLoc));
}
+/// Parsing of simple OpenMP clauses like 'default' or 'proc_bind'.
+///
+/// default-clause:
+/// 'default' '(' 'none' | 'shared' ')
+///
+/// proc_bind-clause:
+/// 'proc_bind' '(' 'master' | 'close' | 'spread' ')
+///
+/// device_type-clause:
+/// 'device_type' '(' 'host' | 'nohost' | 'any' )'
+namespace {
+ struct SimpleClauseData {
+ unsigned Type;
+ SourceLocation Loc;
+ SourceLocation LOpen;
+ SourceLocation TypeLoc;
+ SourceLocation RLoc;
+ SimpleClauseData(unsigned Type, SourceLocation Loc, SourceLocation LOpen,
+ SourceLocation TypeLoc, SourceLocation RLoc)
+ : Type(Type), Loc(Loc), LOpen(LOpen), TypeLoc(TypeLoc), RLoc(RLoc) {}
+ };
+} // anonymous namespace
+
+static Optional<SimpleClauseData>
+parseOpenMPSimpleClause(Parser &P, OpenMPClauseKind Kind) {
+ const Token &Tok = P.getCurToken();
+ SourceLocation Loc = Tok.getLocation();
+ SourceLocation LOpen = P.ConsumeToken();
+ // Parse '('.
+ BalancedDelimiterTracker T(P, tok::l_paren, tok::annot_pragma_openmp_end);
+ if (T.expectAndConsume(diag::err_expected_lparen_after,
+ getOpenMPClauseName(Kind)))
+ return llvm::None;
+
+ unsigned Type = getOpenMPSimpleClauseType(
+ Kind, Tok.isAnnotation() ? "" : P.getPreprocessor().getSpelling(Tok));
+ SourceLocation TypeLoc = Tok.getLocation();
+ if (Tok.isNot(tok::r_paren) && Tok.isNot(tok::comma) &&
+ Tok.isNot(tok::annot_pragma_openmp_end))
+ P.ConsumeAnyToken();
+
+ // Parse ')'.
+ SourceLocation RLoc = Tok.getLocation();
+ if (!T.consumeClose())
+ RLoc = T.getCloseLocation();
+
+ return SimpleClauseData(Type, Loc, LOpen, TypeLoc, RLoc);
+}
+
Parser::DeclGroupPtrTy Parser::ParseOMPDeclareTargetClauses() {
// OpenMP 4.5 syntax with list of entities.
Sema::NamedDeclSetType SameDirectiveDecls;
+ SmallVector<std::tuple<OMPDeclareTargetDeclAttr::MapTypeTy, SourceLocation,
+ NamedDecl *>,
+ 4>
+ DeclareTargetDecls;
+ OMPDeclareTargetDeclAttr::DevTypeTy DT = OMPDeclareTargetDeclAttr::DT_Any;
+ SourceLocation DeviceTypeLoc;
while (Tok.isNot(tok::annot_pragma_openmp_end)) {
OMPDeclareTargetDeclAttr::MapTypeTy MT = OMPDeclareTargetDeclAttr::MT_To;
if (Tok.is(tok::identifier)) {
IdentifierInfo *II = Tok.getIdentifierInfo();
StringRef ClauseName = II->getName();
- // Parse 'to|link' clauses.
- if (!OMPDeclareTargetDeclAttr::ConvertStrToMapTypeTy(ClauseName, MT)) {
- Diag(Tok, diag::err_omp_declare_target_unexpected_clause) << ClauseName;
+ bool IsDeviceTypeClause =
+ getLangOpts().OpenMP >= 50 &&
+ getOpenMPClauseKind(ClauseName) == OMPC_device_type;
+ // Parse 'to|link|device_type' clauses.
+ if (!OMPDeclareTargetDeclAttr::ConvertStrToMapTypeTy(ClauseName, MT) &&
+ !IsDeviceTypeClause) {
+ Diag(Tok, diag::err_omp_declare_target_unexpected_clause)
+ << ClauseName << (getLangOpts().OpenMP >= 50 ? 1 : 0);
break;
}
+ // Parse 'device_type' clause and go to next clause if any.
+ if (IsDeviceTypeClause) {
+ Optional<SimpleClauseData> DevTypeData =
+ parseOpenMPSimpleClause(*this, OMPC_device_type);
+ if (DevTypeData.hasValue()) {
+ if (DeviceTypeLoc.isValid()) {
+ // We already saw another device_type clause, diagnose it.
+ Diag(DevTypeData.getValue().Loc,
+ diag::warn_omp_more_one_device_type_clause);
+ }
+ switch(static_cast<OpenMPDeviceType>(DevTypeData.getValue().Type)) {
+ case OMPC_DEVICE_TYPE_any:
+ DT = OMPDeclareTargetDeclAttr::DT_Any;
+ break;
+ case OMPC_DEVICE_TYPE_host:
+ DT = OMPDeclareTargetDeclAttr::DT_Host;
+ break;
+ case OMPC_DEVICE_TYPE_nohost:
+ DT = OMPDeclareTargetDeclAttr::DT_NoHost;
+ break;
+ case OMPC_DEVICE_TYPE_unknown:
+ llvm_unreachable("Unexpected device_type");
+ }
+ DeviceTypeLoc = DevTypeData.getValue().Loc;
+ }
+ continue;
+ }
ConsumeToken();
}
- auto &&Callback = [this, MT, &SameDirectiveDecls](
- CXXScopeSpec &SS, DeclarationNameInfo NameInfo) {
- Actions.ActOnOpenMPDeclareTargetName(getCurScope(), SS, NameInfo, MT,
- SameDirectiveDecls);
+ auto &&Callback = [this, MT, &DeclareTargetDecls, &SameDirectiveDecls](
+ CXXScopeSpec &SS, DeclarationNameInfo NameInfo) {
+ NamedDecl *ND = Actions.lookupOpenMPDeclareTargetName(
+ getCurScope(), SS, NameInfo, SameDirectiveDecls);
+ if (ND)
+ DeclareTargetDecls.emplace_back(MT, NameInfo.getLoc(), ND);
};
if (ParseOpenMPSimpleVarList(OMPD_declare_target, Callback,
/*AllowScopeSpecifier=*/true))
@@ -812,6 +901,15 @@ Parser::DeclGroupPtrTy Parser::ParseOMPD
}
SkipUntil(tok::annot_pragma_openmp_end, StopBeforeMatch);
ConsumeAnyToken();
+ for (auto &MTLocDecl : DeclareTargetDecls) {
+ OMPDeclareTargetDeclAttr::MapTypeTy MT;
+ SourceLocation Loc;
+ NamedDecl *ND;
+ std::tie(MT, Loc, ND) = MTLocDecl;
+ // device_type clause is applied only to functions.
+ Actions.ActOnOpenMPDeclareTargetName(
+ ND, Loc, MT, isa<VarDecl>(ND) ? OMPDeclareTargetDeclAttr::DT_Any : DT);
+ }
SmallVector<Decl *, 4> Decls(SameDirectiveDecls.begin(),
SameDirectiveDecls.end());
if (Decls.empty())
@@ -1712,6 +1810,7 @@ OMPClause *Parser::ParseOpenMPClause(Ope
case OMPC_allocate:
Clause = ParseOpenMPVarListClause(DKind, CKind, WrongDirective);
break;
+ case OMPC_device_type:
case OMPC_unknown:
Diag(Tok, diag::warn_omp_extra_tokens_at_eol)
<< getOpenMPDirectiveName(DKind);
@@ -1811,29 +1910,12 @@ OMPClause *Parser::ParseOpenMPSingleExpr
///
OMPClause *Parser::ParseOpenMPSimpleClause(OpenMPClauseKind Kind,
bool ParseOnly) {
- SourceLocation Loc = Tok.getLocation();
- SourceLocation LOpen = ConsumeToken();
- // Parse '('.
- BalancedDelimiterTracker T(*this, tok::l_paren, tok::annot_pragma_openmp_end);
- if (T.expectAndConsume(diag::err_expected_lparen_after,
- getOpenMPClauseName(Kind)))
- return nullptr;
-
- unsigned Type = getOpenMPSimpleClauseType(
- Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok));
- SourceLocation TypeLoc = Tok.getLocation();
- if (Tok.isNot(tok::r_paren) && Tok.isNot(tok::comma) &&
- Tok.isNot(tok::annot_pragma_openmp_end))
- ConsumeAnyToken();
-
- // Parse ')'.
- SourceLocation RLoc = Tok.getLocation();
- if (!T.consumeClose())
- RLoc = T.getCloseLocation();
-
- if (ParseOnly)
+ llvm::Optional<SimpleClauseData> Val = parseOpenMPSimpleClause(*this, Kind);
+ if (!Val || ParseOnly)
return nullptr;
- return Actions.ActOnOpenMPSimpleClause(Kind, Type, TypeLoc, LOpen, Loc, RLoc);
+ return Actions.ActOnOpenMPSimpleClause(
+ Kind, Val.getValue().Type, Val.getValue().TypeLoc, Val.getValue().LOpen,
+ Val.getValue().Loc, Val.getValue().RLoc);
}
/// Parsing of OpenMP clauses like 'ordered'.
Modified: cfe/trunk/lib/Sema/Sema.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/Sema.cpp?rev=369775&r1=369774&r2=369775&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/Sema.cpp (original)
+++ cfe/trunk/lib/Sema/Sema.cpp Fri Aug 23 09:11:14 2019
@@ -913,6 +913,10 @@ void Sema::ActOnEndOfTranslationUnitFrag
PerformPendingInstantiations();
}
+ // Finalize analysis of OpenMP-specific constructs.
+ if (LangOpts.OpenMP)
+ finalizeOpenMPDelayedAnalysis();
+
assert(LateParsedInstantiations.empty() &&
"end of TU template instantiation should not create more "
"late-parsed templates");
@@ -1542,8 +1546,9 @@ void Sema::markKnownEmitted(
}
Sema::DeviceDiagBuilder Sema::targetDiag(SourceLocation Loc, unsigned DiagID) {
- if (LangOpts.OpenMP && LangOpts.OpenMPIsDevice)
- return diagIfOpenMPDeviceCode(Loc, DiagID);
+ if (LangOpts.OpenMP)
+ return LangOpts.OpenMPIsDevice ? diagIfOpenMPDeviceCode(Loc, DiagID)
+ : diagIfOpenMPHostCode(Loc, DiagID);
if (getLangOpts().CUDA)
return getLangOpts().CUDAIsDevice ? CUDADiagIfDeviceCode(Loc, DiagID)
: CUDADiagIfHostCode(Loc, DiagID);
Modified: cfe/trunk/lib/Sema/SemaExpr.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaExpr.cpp?rev=369775&r1=369774&r2=369775&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaExpr.cpp (original)
+++ cfe/trunk/lib/Sema/SemaExpr.cpp Fri Aug 23 09:11:14 2019
@@ -15350,9 +15350,13 @@ void Sema::MarkFunctionReferenced(Source
CheckCompleteParameterTypesForMangler(*this, Func, Loc);
Func->markUsed(Context);
+ }
- if (LangOpts.OpenMP && LangOpts.OpenMPIsDevice)
+ if (LangOpts.OpenMP) {
+ if (LangOpts.OpenMPIsDevice)
checkOpenMPDeviceFunction(Loc, Func);
+ else
+ checkOpenMPHostFunction(Loc, Func);
}
}
@@ -17745,4 +17749,4 @@ ExprResult Sema::ActOnObjCAvailabilityCh
return new (Context)
ObjCAvailabilityCheckExpr(Version, AtLoc, RParen, Context.BoolTy);
-}
\ No newline at end of file
+}
Modified: cfe/trunk/lib/Sema/SemaOpenMP.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaOpenMP.cpp?rev=369775&r1=369774&r2=369775&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaOpenMP.cpp (original)
+++ cfe/trunk/lib/Sema/SemaOpenMP.cpp Fri Aug 23 09:11:14 2019
@@ -1556,32 +1556,102 @@ static bool isOpenMPDeviceDelayedContext
!S.isInOpenMPDeclareTargetContext();
}
+namespace {
+/// Status of the function emission on the host/device.
+enum class FunctionEmissionStatus {
+ Emitted,
+ Discarded,
+ Unknown,
+};
+} // anonymous namespace
+
/// Do we know that we will eventually codegen the given function?
-static bool isKnownEmitted(Sema &S, FunctionDecl *FD) {
+static FunctionEmissionStatus isKnownDeviceEmitted(Sema &S, FunctionDecl *FD) {
assert(S.LangOpts.OpenMP && S.LangOpts.OpenMPIsDevice &&
"Expected OpenMP device compilation.");
// Templates are emitted when they're instantiated.
if (FD->isDependentContext())
- return false;
+ return FunctionEmissionStatus::Discarded;
- if (OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(
- FD->getCanonicalDecl()))
- return true;
+ Optional<OMPDeclareTargetDeclAttr::DevTypeTy> DevTy =
+ OMPDeclareTargetDeclAttr::getDeviceType(FD->getCanonicalDecl());
+ if (DevTy.hasValue())
+ return (*DevTy == OMPDeclareTargetDeclAttr::DT_Host)
+ ? FunctionEmissionStatus::Discarded
+ : FunctionEmissionStatus::Emitted;
// Otherwise, the function is known-emitted if it's in our set of
// known-emitted functions.
- return S.DeviceKnownEmittedFns.count(FD) > 0;
+ return (S.DeviceKnownEmittedFns.count(FD) > 0)
+ ? FunctionEmissionStatus::Emitted
+ : FunctionEmissionStatus::Unknown;
}
Sema::DeviceDiagBuilder Sema::diagIfOpenMPDeviceCode(SourceLocation Loc,
unsigned DiagID) {
assert(LangOpts.OpenMP && LangOpts.OpenMPIsDevice &&
"Expected OpenMP device compilation.");
- return DeviceDiagBuilder((isOpenMPDeviceDelayedContext(*this) &&
- !isKnownEmitted(*this, getCurFunctionDecl()))
- ? DeviceDiagBuilder::K_Deferred
- : DeviceDiagBuilder::K_Immediate,
- Loc, DiagID, getCurFunctionDecl(), *this);
+ FunctionEmissionStatus FES =
+ isKnownDeviceEmitted(*this, getCurFunctionDecl());
+ DeviceDiagBuilder::Kind Kind = DeviceDiagBuilder::K_Nop;
+ switch (FES) {
+ case FunctionEmissionStatus::Emitted:
+ Kind = DeviceDiagBuilder::K_Immediate;
+ break;
+ case FunctionEmissionStatus::Unknown:
+ Kind = isOpenMPDeviceDelayedContext(*this) ? DeviceDiagBuilder::K_Deferred
+ : DeviceDiagBuilder::K_Immediate;
+ break;
+ case FunctionEmissionStatus::Discarded:
+ Kind = DeviceDiagBuilder::K_Nop;
+ break;
+ }
+
+ return DeviceDiagBuilder(Kind, Loc, DiagID, getCurFunctionDecl(), *this);
+}
+
+/// Do we know that we will eventually codegen the given function?
+static FunctionEmissionStatus isKnownHostEmitted(Sema &S, FunctionDecl *FD) {
+ assert(S.LangOpts.OpenMP && !S.LangOpts.OpenMPIsDevice &&
+ "Expected OpenMP host compilation.");
+ // In OpenMP 4.5 all the functions are host functions.
+ if (S.LangOpts.OpenMP <= 45)
+ return FunctionEmissionStatus::Emitted;
+
+ Optional<OMPDeclareTargetDeclAttr::DevTypeTy> DevTy =
+ OMPDeclareTargetDeclAttr::getDeviceType(FD->getCanonicalDecl());
+ if (DevTy.hasValue())
+ return (*DevTy == OMPDeclareTargetDeclAttr::DT_NoHost)
+ ? FunctionEmissionStatus::Discarded
+ : FunctionEmissionStatus::Emitted;
+
+ // Otherwise, the function is known-emitted if it's in our set of
+ // known-emitted functions.
+ return (S.DeviceKnownEmittedFns.count(FD) > 0)
+ ? FunctionEmissionStatus::Emitted
+ : FunctionEmissionStatus::Unknown;
+}
+
+Sema::DeviceDiagBuilder Sema::diagIfOpenMPHostCode(SourceLocation Loc,
+ unsigned DiagID) {
+ assert(LangOpts.OpenMP && !LangOpts.OpenMPIsDevice &&
+ "Expected OpenMP host compilation.");
+ FunctionEmissionStatus FES =
+ isKnownHostEmitted(*this, getCurFunctionDecl());
+ DeviceDiagBuilder::Kind Kind = DeviceDiagBuilder::K_Nop;
+ switch (FES) {
+ case FunctionEmissionStatus::Emitted:
+ Kind = DeviceDiagBuilder::K_Immediate;
+ break;
+ case FunctionEmissionStatus::Unknown:
+ Kind = DeviceDiagBuilder::K_Deferred;
+ break;
+ case FunctionEmissionStatus::Discarded:
+ Kind = DeviceDiagBuilder::K_Nop;
+ break;
+ }
+
+ return DeviceDiagBuilder(Kind, Loc, DiagID, getCurFunctionDecl(), *this);
}
void Sema::checkOpenMPDeviceFunction(SourceLocation Loc, FunctionDecl *Callee,
@@ -1589,21 +1659,75 @@ void Sema::checkOpenMPDeviceFunction(Sou
assert(LangOpts.OpenMP && LangOpts.OpenMPIsDevice &&
"Expected OpenMP device compilation.");
assert(Callee && "Callee may not be null.");
+ Callee = Callee->getMostRecentDecl();
FunctionDecl *Caller = getCurFunctionDecl();
+ // host only function are not available on the device.
+ if (Caller &&
+ (isKnownDeviceEmitted(*this, Caller) == FunctionEmissionStatus::Emitted ||
+ (!isOpenMPDeviceDelayedContext(*this) &&
+ isKnownDeviceEmitted(*this, Caller) ==
+ FunctionEmissionStatus::Unknown)) &&
+ isKnownDeviceEmitted(*this, Callee) ==
+ FunctionEmissionStatus::Discarded) {
+ StringRef HostDevTy =
+ getOpenMPSimpleClauseTypeName(OMPC_device_type, OMPC_DEVICE_TYPE_host);
+ Diag(Loc, diag::err_omp_wrong_device_function_call) << HostDevTy << 0;
+ Diag(Callee->getAttr<OMPDeclareTargetDeclAttr>()->getLocation(),
+ diag::note_omp_marked_device_type_here)
+ << HostDevTy;
+ return;
+ }
// If the caller is known-emitted, mark the callee as known-emitted.
// Otherwise, mark the call in our call graph so we can traverse it later.
if ((CheckForDelayedContext && !isOpenMPDeviceDelayedContext(*this)) ||
(!Caller && !CheckForDelayedContext) ||
- (Caller && isKnownEmitted(*this, Caller)))
+ (Caller &&
+ isKnownDeviceEmitted(*this, Caller) == FunctionEmissionStatus::Emitted))
markKnownEmitted(*this, Caller, Callee, Loc,
[CheckForDelayedContext](Sema &S, FunctionDecl *FD) {
- return CheckForDelayedContext && isKnownEmitted(S, FD);
+ return CheckForDelayedContext &&
+ isKnownDeviceEmitted(S, FD) ==
+ FunctionEmissionStatus::Emitted;
});
else if (Caller)
DeviceCallGraph[Caller].insert({Callee, Loc});
}
+void Sema::checkOpenMPHostFunction(SourceLocation Loc, FunctionDecl *Callee,
+ bool CheckCaller) {
+ assert(LangOpts.OpenMP && !LangOpts.OpenMPIsDevice &&
+ "Expected OpenMP host compilation.");
+ assert(Callee && "Callee may not be null.");
+ Callee = Callee->getMostRecentDecl();
+ FunctionDecl *Caller = getCurFunctionDecl();
+
+ // device only function are not available on the host.
+ if (Caller &&
+ isKnownHostEmitted(*this, Caller) == FunctionEmissionStatus::Emitted &&
+ isKnownHostEmitted(*this, Callee) == FunctionEmissionStatus::Discarded) {
+ StringRef NoHostDevTy = getOpenMPSimpleClauseTypeName(
+ OMPC_device_type, OMPC_DEVICE_TYPE_nohost);
+ Diag(Loc, diag::err_omp_wrong_device_function_call) << NoHostDevTy << 1;
+ Diag(Callee->getAttr<OMPDeclareTargetDeclAttr>()->getLocation(),
+ diag::note_omp_marked_device_type_here)
+ << NoHostDevTy;
+ return;
+ }
+ // If the caller is known-emitted, mark the callee as known-emitted.
+ // Otherwise, mark the call in our call graph so we can traverse it later.
+ if ((!CheckCaller && !Caller) ||
+ (Caller &&
+ isKnownHostEmitted(*this, Caller) == FunctionEmissionStatus::Emitted))
+ markKnownEmitted(
+ *this, Caller, Callee, Loc, [CheckCaller](Sema &S, FunctionDecl *FD) {
+ return CheckCaller &&
+ isKnownHostEmitted(S, FD) == FunctionEmissionStatus::Emitted;
+ });
+ else if (Caller)
+ DeviceCallGraph[Caller].insert({Callee, Loc});
+}
+
void Sema::checkOpenMPDeviceExpr(const Expr *E) {
assert(getLangOpts().OpenMP && getLangOpts().OpenMPIsDevice &&
"OpenMP device compilation mode is expected.");
@@ -1970,6 +2094,54 @@ bool Sema::isOpenMPTargetCapturedDecl(co
void Sema::DestroyDataSharingAttributesStack() { delete DSAStack; }
+void Sema::finalizeOpenMPDelayedAnalysis() {
+ assert(LangOpts.OpenMP && "Expected OpenMP compilation mode.");
+ // Diagnose implicit declare target functions and their callees.
+ for (const auto &CallerCallees : DeviceCallGraph) {
+ Optional<OMPDeclareTargetDeclAttr::DevTypeTy> DevTy =
+ OMPDeclareTargetDeclAttr::getDeviceType(
+ CallerCallees.getFirst()->getMostRecentDecl());
+ // Ignore host functions during device analyzis.
+ if (LangOpts.OpenMPIsDevice && DevTy &&
+ *DevTy == OMPDeclareTargetDeclAttr::DT_Host)
+ continue;
+ // Ignore nohost functions during host analyzis.
+ if (!LangOpts.OpenMPIsDevice && DevTy &&
+ *DevTy == OMPDeclareTargetDeclAttr::DT_NoHost)
+ continue;
+ for (const std::pair<CanonicalDeclPtr<FunctionDecl>, SourceLocation>
+ &Callee : CallerCallees.getSecond()) {
+ const FunctionDecl *FD = Callee.first->getMostRecentDecl();
+ Optional<OMPDeclareTargetDeclAttr::DevTypeTy> DevTy =
+ OMPDeclareTargetDeclAttr::getDeviceType(FD);
+ if (LangOpts.OpenMPIsDevice && DevTy &&
+ *DevTy == OMPDeclareTargetDeclAttr::DT_Host) {
+ // Diagnose host function called during device codegen.
+ StringRef HostDevTy = getOpenMPSimpleClauseTypeName(
+ OMPC_device_type, OMPC_DEVICE_TYPE_host);
+ Diag(Callee.second, diag::err_omp_wrong_device_function_call)
+ << HostDevTy << 0;
+ Diag(FD->getAttr<OMPDeclareTargetDeclAttr>()->getLocation(),
+ diag::note_omp_marked_device_type_here)
+ << HostDevTy;
+ continue;
+ }
+ if (!LangOpts.OpenMPIsDevice && DevTy &&
+ *DevTy == OMPDeclareTargetDeclAttr::DT_NoHost) {
+ // Diagnose nohost function called during host codegen.
+ StringRef NoHostDevTy = getOpenMPSimpleClauseTypeName(
+ OMPC_device_type, OMPC_DEVICE_TYPE_nohost);
+ Diag(Callee.second, diag::err_omp_wrong_device_function_call)
+ << NoHostDevTy << 1;
+ Diag(FD->getAttr<OMPDeclareTargetDeclAttr>()->getLocation(),
+ diag::note_omp_marked_device_type_here)
+ << NoHostDevTy;
+ continue;
+ }
+ }
+ }
+}
+
void Sema::StartOpenMPDSABlock(OpenMPDirectiveKind DKind,
const DeclarationNameInfo &DirName,
Scope *CurScope, SourceLocation Loc) {
@@ -4415,6 +4587,7 @@ StmtResult Sema::ActOnOpenMPExecutableDi
case OMPC_reverse_offload:
case OMPC_dynamic_allocators:
case OMPC_atomic_default_mem_order:
+ case OMPC_device_type:
llvm_unreachable("Unexpected clause");
}
for (Stmt *CC : C->children()) {
@@ -9642,6 +9815,7 @@ OMPClause *Sema::ActOnOpenMPSingleExprCl
case OMPC_reverse_offload:
case OMPC_dynamic_allocators:
case OMPC_atomic_default_mem_order:
+ case OMPC_device_type:
llvm_unreachable("Clause is not allowed.");
}
return Res;
@@ -10184,6 +10358,7 @@ static OpenMPDirectiveKind getOpenMPCapt
case OMPC_reverse_offload:
case OMPC_dynamic_allocators:
case OMPC_atomic_default_mem_order:
+ case OMPC_device_type:
llvm_unreachable("Unexpected OpenMP clause.");
}
return CaptureRegion;
@@ -10577,6 +10752,7 @@ OMPClause *Sema::ActOnOpenMPSimpleClause
case OMPC_unified_shared_memory:
case OMPC_reverse_offload:
case OMPC_dynamic_allocators:
+ case OMPC_device_type:
llvm_unreachable("Clause is not allowed.");
}
return Res;
@@ -10755,6 +10931,7 @@ OMPClause *Sema::ActOnOpenMPSingleExprWi
case OMPC_reverse_offload:
case OMPC_dynamic_allocators:
case OMPC_atomic_default_mem_order:
+ case OMPC_device_type:
llvm_unreachable("Clause is not allowed.");
}
return Res;
@@ -10964,6 +11141,7 @@ OMPClause *Sema::ActOnOpenMPClause(OpenM
case OMPC_use_device_ptr:
case OMPC_is_device_ptr:
case OMPC_atomic_default_mem_order:
+ case OMPC_device_type:
llvm_unreachable("Clause is not allowed.");
}
return Res;
@@ -11170,6 +11348,7 @@ OMPClause *Sema::ActOnOpenMPVarListClaus
case OMPC_reverse_offload:
case OMPC_dynamic_allocators:
case OMPC_atomic_default_mem_order:
+ case OMPC_device_type:
llvm_unreachable("Clause is not allowed.");
}
return Res;
@@ -15333,16 +15512,15 @@ void Sema::ActOnFinishOpenMPDeclareTarge
--DeclareTargetNestingLevel;
}
-void Sema::ActOnOpenMPDeclareTargetName(Scope *CurScope,
- CXXScopeSpec &ScopeSpec,
- const DeclarationNameInfo &Id,
- OMPDeclareTargetDeclAttr::MapTypeTy MT,
- NamedDeclSetType &SameDirectiveDecls) {
+NamedDecl *
+Sema::lookupOpenMPDeclareTargetName(Scope *CurScope, CXXScopeSpec &ScopeSpec,
+ const DeclarationNameInfo &Id,
+ NamedDeclSetType &SameDirectiveDecls) {
LookupResult Lookup(*this, Id, LookupOrdinaryName);
LookupParsedName(Lookup, CurScope, &ScopeSpec, true);
if (Lookup.isAmbiguous())
- return;
+ return nullptr;
Lookup.suppressDiagnostics();
if (!Lookup.isSingleResult()) {
@@ -15353,33 +15531,56 @@ void Sema::ActOnOpenMPDeclareTargetName(
diagnoseTypo(Corrected, PDiag(diag::err_undeclared_var_use_suggest)
<< Id.getName());
checkDeclIsAllowedInOpenMPTarget(nullptr, Corrected.getCorrectionDecl());
- return;
+ return nullptr;
}
Diag(Id.getLoc(), diag::err_undeclared_var_use) << Id.getName();
- return;
+ return nullptr;
}
NamedDecl *ND = Lookup.getAsSingle<NamedDecl>();
- if (isa<VarDecl>(ND) || isa<FunctionDecl>(ND) ||
- isa<FunctionTemplateDecl>(ND)) {
- if (!SameDirectiveDecls.insert(cast<NamedDecl>(ND->getCanonicalDecl())))
- Diag(Id.getLoc(), diag::err_omp_declare_target_multiple) << Id.getName();
- llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
- OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(
- cast<ValueDecl>(ND));
- if (!Res) {
- auto *A = OMPDeclareTargetDeclAttr::CreateImplicit(Context, MT);
- ND->addAttr(A);
- if (ASTMutationListener *ML = Context.getASTMutationListener())
- ML->DeclarationMarkedOpenMPDeclareTarget(ND, A);
- checkDeclIsAllowedInOpenMPTarget(nullptr, ND, Id.getLoc());
- } else if (*Res != MT) {
- Diag(Id.getLoc(), diag::err_omp_declare_target_to_and_link)
- << Id.getName();
- }
- } else {
+ if (!isa<VarDecl>(ND) && !isa<FunctionDecl>(ND) &&
+ !isa<FunctionTemplateDecl>(ND)) {
Diag(Id.getLoc(), diag::err_omp_invalid_target_decl) << Id.getName();
+ return nullptr;
+ }
+ if (!SameDirectiveDecls.insert(cast<NamedDecl>(ND->getCanonicalDecl())))
+ Diag(Id.getLoc(), diag::err_omp_declare_target_multiple) << Id.getName();
+ return ND;
+}
+
+void Sema::ActOnOpenMPDeclareTargetName(
+ NamedDecl *ND, SourceLocation Loc, OMPDeclareTargetDeclAttr::MapTypeTy MT,
+ OMPDeclareTargetDeclAttr::DevTypeTy DT) {
+ assert((isa<VarDecl>(ND) || isa<FunctionDecl>(ND) ||
+ isa<FunctionTemplateDecl>(ND)) &&
+ "Expected variable, function or function template.");
+
+ // Diagnose marking after use as it may lead to incorrect diagnosis and
+ // codegen.
+ if (LangOpts.OpenMP >= 50 &&
+ (ND->isUsed(/*CheckUsedAttr=*/false) || ND->isReferenced()))
+ Diag(Loc, diag::warn_omp_declare_target_after_first_use);
+
+ Optional<OMPDeclareTargetDeclAttr::DevTypeTy> DevTy =
+ OMPDeclareTargetDeclAttr::getDeviceType(cast<ValueDecl>(ND));
+ if (DevTy.hasValue() && *DevTy != DT) {
+ Diag(Loc, diag::err_omp_device_type_mismatch)
+ << OMPDeclareTargetDeclAttr::ConvertDevTypeTyToStr(DT)
+ << OMPDeclareTargetDeclAttr::ConvertDevTypeTyToStr(*DevTy);
+ return;
+ }
+ Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
+ OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(cast<ValueDecl>(ND));
+ if (!Res) {
+ auto *A = OMPDeclareTargetDeclAttr::CreateImplicit(Context, MT, DT,
+ SourceRange(Loc, Loc));
+ ND->addAttr(A);
+ if (ASTMutationListener *ML = Context.getASTMutationListener())
+ ML->DeclarationMarkedOpenMPDeclareTarget(ND, A);
+ checkDeclIsAllowedInOpenMPTarget(nullptr, ND, Loc);
+ } else if (*Res != MT) {
+ Diag(Loc, diag::err_omp_declare_target_to_and_link) << ND;
}
}
@@ -15453,8 +15654,14 @@ void Sema::checkDeclIsAllowedInOpenMPTar
return;
}
// Mark the function as must be emitted for the device.
- if (LangOpts.OpenMPIsDevice && Res.hasValue() && IdLoc.isValid())
+ Optional<OMPDeclareTargetDeclAttr::DevTypeTy> DevTy =
+ OMPDeclareTargetDeclAttr::getDeviceType(FD);
+ if (LangOpts.OpenMPIsDevice && Res.hasValue() && IdLoc.isValid() &&
+ *DevTy != OMPDeclareTargetDeclAttr::DT_Host)
checkOpenMPDeviceFunction(IdLoc, FD, /*CheckForDelayedContext=*/false);
+ if (!LangOpts.OpenMPIsDevice && Res.hasValue() && IdLoc.isValid() &&
+ *DevTy != OMPDeclareTargetDeclAttr::DT_NoHost)
+ checkOpenMPHostFunction(IdLoc, FD, /*CheckCaller=*/false);
}
if (auto *VD = dyn_cast<ValueDecl>(D)) {
// Problem if any with var declared with incomplete type will be reported
@@ -15467,7 +15674,8 @@ void Sema::checkDeclIsAllowedInOpenMPTar
if (isa<VarDecl>(D) || isa<FunctionDecl>(D) ||
isa<FunctionTemplateDecl>(D)) {
auto *A = OMPDeclareTargetDeclAttr::CreateImplicit(
- Context, OMPDeclareTargetDeclAttr::MT_To);
+ Context, OMPDeclareTargetDeclAttr::MT_To,
+ OMPDeclareTargetDeclAttr::DT_Any, SourceRange(IdLoc, IdLoc));
D->addAttr(A);
if (ASTMutationListener *ML = Context.getASTMutationListener())
ML->DeclarationMarkedOpenMPDeclareTarget(D, A);
Modified: cfe/trunk/lib/Serialization/ASTReaderDecl.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Serialization/ASTReaderDecl.cpp?rev=369775&r1=369774&r2=369775&view=diff
==============================================================================
--- cfe/trunk/lib/Serialization/ASTReaderDecl.cpp (original)
+++ cfe/trunk/lib/Serialization/ASTReaderDecl.cpp Fri Aug 23 09:11:14 2019
@@ -4571,12 +4571,15 @@ void ASTDeclReader::UpdateDecl(Decl *D,
break;
}
- case UPD_DECL_MARKED_OPENMP_DECLARETARGET:
+ case UPD_DECL_MARKED_OPENMP_DECLARETARGET: {
+ OMPDeclareTargetDeclAttr::MapTypeTy MapType =
+ static_cast<OMPDeclareTargetDeclAttr::MapTypeTy>(Record.readInt());
+ OMPDeclareTargetDeclAttr::DevTypeTy DevType =
+ static_cast<OMPDeclareTargetDeclAttr::DevTypeTy>(Record.readInt());
D->addAttr(OMPDeclareTargetDeclAttr::CreateImplicit(
- Reader.getContext(),
- static_cast<OMPDeclareTargetDeclAttr::MapTypeTy>(Record.readInt()),
- ReadSourceRange()));
+ Reader.getContext(), MapType, DevType, ReadSourceRange()));
break;
+ }
case UPD_ADDED_ATTR_TO_RECORD:
AttrVec Attrs;
Modified: cfe/trunk/test/OpenMP/declare_target_ast_print.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/declare_target_ast_print.cpp?rev=369775&r1=369774&r2=369775&view=diff
==============================================================================
--- cfe/trunk/test/OpenMP/declare_target_ast_print.cpp (original)
+++ cfe/trunk/test/OpenMP/declare_target_ast_print.cpp Fri Aug 23 09:11:14 2019
@@ -2,6 +2,10 @@
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -I %S/Inputs -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp -std=c++11 -include-pch %t -fsyntax-only -I %S/Inputs -verify %s -ast-print | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 -I %S/Inputs -ast-print %s | FileCheck %s --check-prefix=CHECK --check-prefix=OMP50
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -I %S/Inputs -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -I %S/Inputs -verify %s -ast-print | FileCheck %s --check-prefix=CHECK --check-prefix=OMP50
+
// RUN: %clang_cc1 -verify -fopenmp-simd -I %S/Inputs -ast-print %s | FileCheck %s
// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -I %S/Inputs -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp-simd -std=c++11 -include-pch %t -fsyntax-only -I %S/Inputs -verify %s -ast-print | FileCheck %s
@@ -10,7 +14,29 @@
#ifndef HEADER
#define HEADER
+#if _OPENMP == 201811
+void bar();
+#pragma omp declare target to(bar) device_type(any)
+// OMP50: #pragma omp declare target{{$}}
+// OMP50: void bar();
+// OMP50: #pragma omp end declare target{{$}}
+void baz();
+#pragma omp declare target to(baz) device_type(nohost)
+// OMP50: #pragma omp declare target device_type(nohost){{$}}
+// OMP50: void baz();
+// OMP50: #pragma omp end declare target{{$}}
+void bazz();
+#pragma omp declare target to(bazz) device_type(host)
+// OMP50: #pragma omp declare target device_type(host){{$}}
+// OMP50: void bazz();
+// OMP50: #pragma omp end declare target{{$}}
+#endif // _OPENMP
+
int out_decl_target = 0;
+#if _OPENMP == 201811
+#pragma omp declare target (out_decl_target)
+#endif // _OPENMP
+
// CHECK: #pragma omp declare target{{$}}
// CHECK: int out_decl_target = 0;
// CHECK: #pragma omp end declare target{{$}}
Modified: cfe/trunk/test/OpenMP/declare_target_codegen.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/declare_target_codegen.cpp?rev=369775&r1=369774&r2=369775&view=diff
==============================================================================
--- cfe/trunk/test/OpenMP/declare_target_codegen.cpp (original)
+++ cfe/trunk/test/OpenMP/declare_target_codegen.cpp Fri Aug 23 09:11:14 2019
@@ -3,6 +3,14 @@
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -o - | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - -fopenmp-version=50 -DOMP5 | FileCheck %s --check-prefix HOST5
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc -fopenmp-version=50 -DOMP5
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -fopenmp-version=50 -DOMP5 | FileCheck %s --check-prefix DEV5
+
+// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - -fopenmp-version=50 -DOMP5 | FileCheck %s --check-prefix SIMD-ONLY
+// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc -fopenmp-version=50 -DOMP5
+// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -fopenmp-version=50 -DOMP5 | FileCheck %s --check-prefix SIMD-ONLY
+
// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o -| FileCheck %s --check-prefix SIMD-ONLY
// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t
@@ -32,13 +40,13 @@
// CHECK-DAG: [[STAT:@.+stat]] = internal global %struct.S zeroinitializer,
// CHECK-DAG: [[STAT_REF:@.+]] = internal constant %struct.S* [[STAT]]
// CHECK-DAG: @out_decl_target ={{ dso_local | }}global i32 0,
-// CHECK-DAG: @llvm.used = appending global [2 x i8*] [i8* bitcast (void ()* @__omp_offloading__{{.+}}_globals_l[[@LINE+80]]_ctor to i8*), i8* bitcast (void ()* @__omp_offloading__{{.+}}_stat_l[[@LINE+81]]_ctor to i8*)],
+// CHECK-DAG: @llvm.used = appending global [2 x i8*] [i8* bitcast (void ()* @__omp_offloading__{{.+}}_globals_l[[@LINE+84]]_ctor to i8*), i8* bitcast (void ()* @__omp_offloading__{{.+}}_stat_l[[@LINE+85]]_ctor to i8*)],
// CHECK-DAG: @llvm.compiler.used = appending global [1 x i8*] [i8* bitcast (%struct.S** [[STAT_REF]] to i8*)],
// CHECK-DAG: define {{.*}}i32 @{{.*}}{{foo|bar|baz2|baz3|FA|f_method}}{{.*}}()
// CHECK-DAG: define {{.*}}void @{{.*}}TemplateClass{{.*}}(%class.TemplateClass* %{{.*}})
// CHECK-DAG: define {{.*}}i32 @{{.*}}TemplateClass{{.*}}f_method{{.*}}(%class.TemplateClass* %{{.*}})
-// CHECK-DAG: define {{.*}}void @__omp_offloading__{{.*}}_globals_l[[@LINE+74]]_ctor()
+// CHECK-DAG: define {{.*}}void @__omp_offloading__{{.*}}_globals_l[[@LINE+78]]_ctor()
#ifndef HEADER
#define HEADER
@@ -68,6 +76,10 @@ int hhh = 0;
#pragma omp declare target link(eee, fff, ggg, hhh)
int out_decl_target = 0;
+#ifdef OMP5
+#pragma omp declare target(out_decl_target)
+#endif
+
#pragma omp declare target
void lambda () {
#ifdef __cpp_lambdas
@@ -224,4 +236,18 @@ int main() {
// CHECK-DAG: !{i32 1, !"ccc", i32 0, i32 {{[0-9]+}}}
// CHECK-DAG: !{{{.+}}virtual_foo
+#ifdef OMP5
+void host_fun() {}
+#pragma omp declare target to(host_fun) device_type(host)
+void device_fun() {}
+#pragma omp declare target to(device_fun) device_type(nohost)
+// HOST5-NOT: define {{.*}}void {{.*}}device_fun{{.*}}
+// HOST5: define {{.*}}void {{.*}}host_fun{{.*}}
+// HOST5-NOT: define {{.*}}void {{.*}}device_fun{{.*}}
+
+// DEV5-NOT: define {{.*}}void {{.*}}host_fun{{.*}}
+// DEV5: define {{.*}}void {{.*}}device_fun{{.*}}
+// DEV5-NOT: define {{.*}}void {{.*}}host_fun{{.*}}
+#endif // OMP5
+
#endif // HEADER
Modified: cfe/trunk/test/OpenMP/declare_target_messages.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/declare_target_messages.cpp?rev=369775&r1=369774&r2=369775&view=diff
==============================================================================
--- cfe/trunk/test/OpenMP/declare_target_messages.cpp (original)
+++ cfe/trunk/test/OpenMP/declare_target_messages.cpp Fri Aug 23 09:11:14 2019
@@ -1,7 +1,10 @@
-// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp -fnoopenmp-use-tls -ferror-limit 100 -o - %s
-// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp5 -fopenmp -fopenmp-version=50 -fnoopenmp-use-tls -ferror-limit 100 -o - %s
-
-// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp-simd -fnoopenmp-use-tls -ferror-limit 100 -o - %s
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp45 -fopenmp -fnoopenmp-use-tls -ferror-limit 100 -o - %s
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp5,host5 -fopenmp -fopenmp-version=50 -fnoopenmp-use-tls -ferror-limit 100 -o - %s
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp5,dev5 -fopenmp -fopenmp-is-device -fopenmp-targets=x86_64-apple-macos10.7.0 -aux-triple x86_64-apple-macos10.7.0 -fopenmp-version=50 -fnoopenmp-use-tls -ferror-limit 100 -o - %s
+
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp5,host5 -fopenmp-simd -fopenmp-version=50 -fnoopenmp-use-tls -ferror-limit 100 -o - %s
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp5,host5 -fopenmp-simd -fopenmp-is-device -fopenmp-version=50 -fnoopenmp-use-tls -ferror-limit 100 -o - %s
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp45 -fopenmp-simd -fnoopenmp-use-tls -ferror-limit 100 -o - %s
#pragma omp end declare target // expected-error {{unexpected OpenMP directive '#pragma omp end declare target'}}
@@ -14,17 +17,23 @@ __thread int t; // expected-note {{defin
void f();
#pragma omp end declare target shared(a) // expected-warning {{extra tokens at the end of '#pragma omp end declare target' are ignored}}
-#pragma omp declare target map(a) // expected-error {{unexpected 'map' clause, only 'to' or 'link' clauses expected}}
+#pragma omp declare target map(a) // omp45-error {{unexpected 'map' clause, only 'to' or 'link' clauses expected}} omp5-error {{unexpected 'map' clause, only 'to', 'link' or 'device_type' clauses expected}}
#pragma omp declare target to(foo1) // expected-error {{use of undeclared identifier 'foo1'}}
#pragma omp declare target link(foo2) // expected-error {{use of undeclared identifier 'foo2'}}
+#pragma omp declare target to(f) device_type(any) device_type(any) device_type(host) // omp45-error {{unexpected 'device_type' clause, only 'to' or 'link' clauses expected}} omp5-warning 2 {{more than one 'device_type' clause is specified}} omp5-error {{'device_type(host)' does not match previously specified 'device_type(any)' for the same declaration}}
+
void c();
void func() {} // expected-note {{'func' defined here}}
-#pragma omp declare target link(func) allocate(a) // expected-error {{function name is not allowed in 'link' clause}} expected-error {{unexpected 'allocate' clause, only 'to' or 'link' clauses expected}}
+#pragma omp declare target link(func) allocate(a) // expected-error {{function name is not allowed in 'link' clause}} omp45-error {{unexpected 'allocate' clause, only 'to' or 'link' clauses expected}} omp5-error {{unexpected 'allocate' clause, only 'to', 'link' or 'device_type' clauses expected}}
+
+void bar();
+void baz() {bar();}
+#pragma omp declare target(bar) // omp5-warning {{declaration marked as declare target after first use, it may lead to incorrect results}}
extern int b;
@@ -152,4 +161,30 @@ namespace {
#pragma omp declare target to(x) to(x) // expected-error {{'x' appears multiple times in clauses on the same declare target directive}}
#pragma omp declare target link(x) // expected-error {{'x' must not appear in both clauses 'to' and 'link'}}
+void bazz() {}
+#pragma omp declare target to(bazz) device_type(nohost) // omp45-error {{unexpected 'device_type' clause, only 'to' or 'link' clauses expected}}
+void bazzz() {bazz();}
+#pragma omp declare target to(bazzz) device_type(nohost) // omp45-error {{unexpected 'device_type' clause, only 'to' or 'link' clauses expected}}
+void any() {bazz();}
+void host1() {bazz();}
+#pragma omp declare target to(host1) device_type(host) // omp45-error {{unexpected 'device_type' clause, only 'to' or 'link' clauses expected}} dev5-note 2 {{marked as 'device_type(host)' here}}
+void host2() {bazz();}
+#pragma omp declare target to(host2)
+void device() {host1();}
+#pragma omp declare target to(device) device_type(nohost) // omp45-error {{unexpected 'device_type' clause, only 'to' or 'link' clauses expected}} host5-note 2 {{marked as 'device_type(nohost)' here}}
+void host3() {host1();}
+#pragma omp declare target to(host3)
+
+#pragma omp declare target
+void any1() {any();}
+void any2() {host1();} // dev5-error {{function with 'device_type(host)' is not available on device}}
+void any3() {device();} // host5-error {{function with 'device_type(nohost)' is not available on host}}
+void any4() {any2();}
+#pragma omp end declare target
+
+void any5() {any();}
+void any6() {host1();} // dev5-error {{function with 'device_type(host)' is not available on device}}
+void any7() {device();} // host5-error {{function with 'device_type(nohost)' is not available on host}}
+void any8() {any2();}
+
#pragma omp declare target // expected-error {{expected '#pragma omp end declare target'}} expected-note {{to match this '#pragma omp declare target'}}
Modified: cfe/trunk/test/OpenMP/nvptx_allocate_messages.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/nvptx_allocate_messages.cpp?rev=369775&r1=369774&r2=369775&view=diff
==============================================================================
--- cfe/trunk/test/OpenMP/nvptx_allocate_messages.cpp (original)
+++ cfe/trunk/test/OpenMP/nvptx_allocate_messages.cpp Fri Aug 23 09:11:14 2019
@@ -58,7 +58,7 @@ template <class T> T foo() {
#pragma omp allocate(v) allocator(omp_cgroup_mem_alloc)
v = ST<T>::m;
#if defined(DEVICE) && !defined(REQUIRES)
-// expected-error at +2 2 {{expected an allocator expression inside of the target region; provide an allocator expression or use 'requires' directive with the 'dynamic_allocators' clause}}
+// expected-error at +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;
Modified: cfe/trunk/test/OpenMP/nvptx_asm_delayed_diags.c
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/nvptx_asm_delayed_diags.c?rev=369775&r1=369774&r2=369775&view=diff
==============================================================================
--- cfe/trunk/test/OpenMP/nvptx_asm_delayed_diags.c (original)
+++ cfe/trunk/test/OpenMP/nvptx_asm_delayed_diags.c Fri Aug 23 09:11:14 2019
@@ -2,6 +2,10 @@
// RUN: %clang_cc1 -verify -fopenmp -x c -triple nvptx-unknown-unknown -aux-triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only -Wuninitialized
// RUN: %clang_cc1 -verify -DDIAGS -DIMMEDIATE -fopenmp -x c -triple nvptx-unknown-unknown -aux-triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only -Wuninitialized
// RUN: %clang_cc1 -verify -DDIAGS -DDELAYED -fopenmp -x c -triple nvptx-unknown-unknown -aux-triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only -Wuninitialized
+// RUN: %clang_cc1 -fopenmp -x c -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -fopenmp-version=50 -emit-llvm-bc %s -o %t-x86-host.bc
+// RUN: %clang_cc1 -verify=expected,omp5 -fopenmp -x c -triple nvptx-unknown-unknown -aux-triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -fopenmp-version=50 %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,omp5 -DDIAGS -DOMP5 -DIMMEDIATE -fopenmp -fopenmp-version=50 -x c -triple nvptx-unknown-unknown -aux-triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,omp5 -DDIAGS -DOMP5 -DDELAYED -fopenmp -fopenmp-version=50 -x c -triple nvptx-unknown-unknown -aux-triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only -Wuninitialized
// REQUIRES: x86-registered-target
// REQUIRES: nvptx-registered-target
@@ -9,6 +13,22 @@
// expected-no-diagnostics
#endif // DIAGS
+#ifdef OMP5
+void bar(int r) {
+#ifdef IMMEDIATE
+// omp5-error at +4 {{invalid input constraint 'mx' in asm}}
+#endif // IMMEDIATE
+ __asm__("PR3908 %[lf] %[xx] %[li] %[r]"
+ : [ r ] "+r"(r)
+ : [ lf ] "mx"(0), [ li ] "mr"(0), [ xx ] "x"((double)(0)));
+}
+#ifdef IMMEDIATE
+#pragma omp declare target to(bar) device_type(nohost)
+#else
+#pragma omp declare target to(bar) device_type(host)
+#endif // IMMEDIATE
+#endif // OMP5
+
void foo(int r) {
#ifdef IMMEDIATE
// expected-error at +4 {{invalid input constraint 'mx' in asm}}
Modified: cfe/trunk/test/OpenMP/target_vla_messages.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/target_vla_messages.cpp?rev=369775&r1=369774&r2=369775&view=diff
==============================================================================
--- cfe/trunk/test/OpenMP/target_vla_messages.cpp (original)
+++ cfe/trunk/test/OpenMP/target_vla_messages.cpp Fri Aug 23 09:11:14 2019
@@ -47,7 +47,7 @@ void target_template(int arg) {
#pragma omp target
{
#ifdef NO_VLA
- // expected-error at +2 2 {{variable length arrays are not supported for the current target}}
+ // expected-error at +2 {{variable length arrays are not supported for the current target}}
#endif
T vla[arg];
}
More information about the cfe-commits
mailing list