aboutsummaryrefslogtreecommitdiff
path: root/clang/lib
diff options
context:
space:
mode:
Diffstat (limited to 'clang/lib')
-rw-r--r--clang/lib/AST/ASTContext.cpp17
-rw-r--r--clang/lib/AST/Decl.cpp8
-rw-r--r--clang/lib/AST/ItaniumMangle.cpp29
-rw-r--r--clang/lib/AST/OpenMPClause.cpp2
-rw-r--r--clang/lib/Basic/Targets/NVPTX.cpp2
-rw-r--r--clang/lib/CodeGen/CGBuiltin.cpp2
-rw-r--r--clang/lib/CodeGen/CGDecl.cpp2
-rw-r--r--clang/lib/CodeGen/CGExpr.cpp2
-rw-r--r--clang/lib/CodeGen/CGOpenMPRuntime.cpp51
-rw-r--r--clang/lib/CodeGen/CGOpenMPRuntime.h2
-rw-r--r--clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp6
-rw-r--r--clang/lib/CodeGen/CGOpenMPRuntimeGPU.h2
-rw-r--r--clang/lib/CodeGen/CGStmtOpenMP.cpp4
-rw-r--r--clang/lib/CodeGen/CGVTables.cpp2
-rw-r--r--clang/lib/CodeGen/CodeGenModule.cpp12
-rw-r--r--clang/lib/CodeGen/Targets/NVPTX.cpp3
-rw-r--r--clang/lib/Driver/ToolChains/Clang.cpp8
-rw-r--r--clang/lib/Driver/ToolChains/Flang.cpp7
-rw-r--r--clang/lib/Frontend/CompilerInstance.cpp2
-rw-r--r--clang/lib/Frontend/CompilerInvocation.cpp30
-rw-r--r--clang/lib/Frontend/InitPreprocessor.cpp3
-rw-r--r--clang/lib/Lex/LiteralSupport.cpp2
-rw-r--r--clang/lib/Sema/Sema.cpp8
-rw-r--r--clang/lib/Sema/SemaDecl.cpp15
-rw-r--r--clang/lib/Sema/SemaDeclCXX.cpp2
-rw-r--r--clang/lib/Sema/SemaExpr.cpp5
-rw-r--r--clang/lib/Sema/SemaOpenMP.cpp20
-rw-r--r--clang/lib/Sema/SemaType.cpp10
28 files changed, 136 insertions, 122 deletions
diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp
index 65132b9..2b84673 100644
--- a/clang/lib/AST/ASTContext.cpp
+++ b/clang/lib/AST/ASTContext.cpp
@@ -1668,11 +1668,11 @@ const llvm::fltSemantics &ASTContext::getFloatTypeSemantics(QualType T) const {
case BuiltinType::Ibm128:
return Target->getIbm128Format();
case BuiltinType::LongDouble:
- if (getLangOpts().OpenMP && getLangOpts().OpenMPIsDevice)
+ if (getLangOpts().OpenMP && getLangOpts().OpenMPIsTargetDevice)
return AuxTarget->getLongDoubleFormat();
return Target->getLongDoubleFormat();
case BuiltinType::Float128:
- if (getLangOpts().OpenMP && getLangOpts().OpenMPIsDevice)
+ if (getLangOpts().OpenMP && getLangOpts().OpenMPIsTargetDevice)
return AuxTarget->getFloat128Format();
return Target->getFloat128Format();
}
@@ -2118,7 +2118,8 @@ TypeInfo ASTContext::getTypeInfoImpl(const Type *T) const {
Width = Target->getBFloat16Width();
Align = Target->getBFloat16Align();
} else if ((getLangOpts().SYCLIsDevice ||
- (getLangOpts().OpenMP && getLangOpts().OpenMPIsDevice)) &&
+ (getLangOpts().OpenMP &&
+ getLangOpts().OpenMPIsTargetDevice)) &&
AuxTarget->hasBFloat16Type()) {
Width = AuxTarget->getBFloat16Width();
Align = AuxTarget->getBFloat16Align();
@@ -2127,11 +2128,11 @@ TypeInfo ASTContext::getTypeInfoImpl(const Type *T) const {
case BuiltinType::Float16:
case BuiltinType::Half:
if (Target->hasFloat16Type() || !getLangOpts().OpenMP ||
- !getLangOpts().OpenMPIsDevice) {
+ !getLangOpts().OpenMPIsTargetDevice) {
Width = Target->getHalfWidth();
Align = Target->getHalfAlign();
} else {
- assert(getLangOpts().OpenMP && getLangOpts().OpenMPIsDevice &&
+ assert(getLangOpts().OpenMP && getLangOpts().OpenMPIsTargetDevice &&
"Expected OpenMP device compilation.");
Width = AuxTarget->getHalfWidth();
Align = AuxTarget->getHalfAlign();
@@ -2150,7 +2151,7 @@ TypeInfo ASTContext::getTypeInfoImpl(const Type *T) const {
Align = Target->getIbm128Align();
break;
case BuiltinType::LongDouble:
- if (getLangOpts().OpenMP && getLangOpts().OpenMPIsDevice &&
+ if (getLangOpts().OpenMP && getLangOpts().OpenMPIsTargetDevice &&
(Target->getLongDoubleWidth() != AuxTarget->getLongDoubleWidth() ||
Target->getLongDoubleAlign() != AuxTarget->getLongDoubleAlign())) {
Width = AuxTarget->getLongDoubleWidth();
@@ -2162,11 +2163,11 @@ TypeInfo ASTContext::getTypeInfoImpl(const Type *T) const {
break;
case BuiltinType::Float128:
if (Target->hasFloat128Type() || !getLangOpts().OpenMP ||
- !getLangOpts().OpenMPIsDevice) {
+ !getLangOpts().OpenMPIsTargetDevice) {
Width = Target->getFloat128Width();
Align = Target->getFloat128Align();
} else {
- assert(getLangOpts().OpenMP && getLangOpts().OpenMPIsDevice &&
+ assert(getLangOpts().OpenMP && getLangOpts().OpenMPIsTargetDevice &&
"Expected OpenMP device compilation.");
Width = AuxTarget->getFloat128Width();
Align = AuxTarget->getFloat128Align();
diff --git a/clang/lib/AST/Decl.cpp b/clang/lib/AST/Decl.cpp
index 6e88bf8..fbc45fb 100644
--- a/clang/lib/AST/Decl.cpp
+++ b/clang/lib/AST/Decl.cpp
@@ -817,7 +817,8 @@ LinkageComputer::getLVForNamespaceScopeDecl(const NamedDecl *D,
// OpenMP target declare device functions are not callable from the host so
// they should not be exported from the device image. This applies to all
// functions as the host-callable kernel functions are emitted at codegen.
- if (Context.getLangOpts().OpenMP && Context.getLangOpts().OpenMPIsDevice &&
+ if (Context.getLangOpts().OpenMP &&
+ Context.getLangOpts().OpenMPIsTargetDevice &&
((Context.getTargetInfo().getTriple().isAMDGPU() ||
Context.getTargetInfo().getTriple().isNVPTX()) ||
OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(Function)))
@@ -1005,7 +1006,8 @@ LinkageComputer::getLVForClassMember(const NamedDecl *D,
// they should not be exported from the device image. This applies to all
// functions as the host-callable kernel functions are emitted at codegen.
ASTContext &Context = D->getASTContext();
- if (Context.getLangOpts().OpenMP && Context.getLangOpts().OpenMPIsDevice &&
+ if (Context.getLangOpts().OpenMP &&
+ Context.getLangOpts().OpenMPIsTargetDevice &&
((Context.getTargetInfo().getTriple().isAMDGPU() ||
Context.getTargetInfo().getTriple().isNVPTX()) ||
OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(MD)))
@@ -3564,7 +3566,7 @@ unsigned FunctionDecl::getBuiltinID(bool ConsiderWrapperFunctions) const {
// library, none of the predefined library functions except printf and malloc
// should be treated as a builtin i.e. 0 should be returned for them.
if (Context.getTargetInfo().getTriple().isAMDGCN() &&
- Context.getLangOpts().OpenMPIsDevice &&
+ Context.getLangOpts().OpenMPIsTargetDevice &&
Context.BuiltinInfo.isPredefinedLibFunction(BuiltinID) &&
!(BuiltinID == Builtin::BIprintf || BuiltinID == Builtin::BImalloc))
return 0;
diff --git a/clang/lib/AST/ItaniumMangle.cpp b/clang/lib/AST/ItaniumMangle.cpp
index fb5bdfa..5ed76e1 100644
--- a/clang/lib/AST/ItaniumMangle.cpp
+++ b/clang/lib/AST/ItaniumMangle.cpp
@@ -3124,27 +3124,30 @@ void CXXNameMangler::mangleType(const BuiltinType *T) {
Out << 'd';
break;
case BuiltinType::LongDouble: {
- const TargetInfo *TI = getASTContext().getLangOpts().OpenMP &&
- getASTContext().getLangOpts().OpenMPIsDevice
- ? getASTContext().getAuxTargetInfo()
- : &getASTContext().getTargetInfo();
+ const TargetInfo *TI =
+ getASTContext().getLangOpts().OpenMP &&
+ getASTContext().getLangOpts().OpenMPIsTargetDevice
+ ? getASTContext().getAuxTargetInfo()
+ : &getASTContext().getTargetInfo();
Out << TI->getLongDoubleMangling();
break;
}
case BuiltinType::Float128: {
- const TargetInfo *TI = getASTContext().getLangOpts().OpenMP &&
- getASTContext().getLangOpts().OpenMPIsDevice
- ? getASTContext().getAuxTargetInfo()
- : &getASTContext().getTargetInfo();
+ const TargetInfo *TI =
+ getASTContext().getLangOpts().OpenMP &&
+ getASTContext().getLangOpts().OpenMPIsTargetDevice
+ ? getASTContext().getAuxTargetInfo()
+ : &getASTContext().getTargetInfo();
Out << TI->getFloat128Mangling();
break;
}
case BuiltinType::BFloat16: {
- const TargetInfo *TI = ((getASTContext().getLangOpts().OpenMP &&
- getASTContext().getLangOpts().OpenMPIsDevice) ||
- getASTContext().getLangOpts().SYCLIsDevice)
- ? getASTContext().getAuxTargetInfo()
- : &getASTContext().getTargetInfo();
+ const TargetInfo *TI =
+ ((getASTContext().getLangOpts().OpenMP &&
+ getASTContext().getLangOpts().OpenMPIsTargetDevice) ||
+ getASTContext().getLangOpts().SYCLIsDevice)
+ ? getASTContext().getAuxTargetInfo()
+ : &getASTContext().getTargetInfo();
Out << TI->getBFloat16Mangling();
break;
}
diff --git a/clang/lib/AST/OpenMPClause.cpp b/clang/lib/AST/OpenMPClause.cpp
index 98870b3..4c89582 100644
--- a/clang/lib/AST/OpenMPClause.cpp
+++ b/clang/lib/AST/OpenMPClause.cpp
@@ -2712,7 +2712,7 @@ TargetOMPContext::TargetOMPContext(
ASTContext &ASTCtx, std::function<void(StringRef)> &&DiagUnknownTrait,
const FunctionDecl *CurrentFunctionDecl,
ArrayRef<llvm::omp::TraitProperty> ConstructTraits)
- : OMPContext(ASTCtx.getLangOpts().OpenMPIsDevice,
+ : OMPContext(ASTCtx.getLangOpts().OpenMPIsTargetDevice,
ASTCtx.getTargetInfo().getTriple()),
FeatureValidityCheck([&](StringRef FeatureName) {
return ASTCtx.getTargetInfo().isValidFeatureName(FeatureName);
diff --git a/clang/lib/Basic/Targets/NVPTX.cpp b/clang/lib/Basic/Targets/NVPTX.cpp
index cfcf4ca..b83bb65 100644
--- a/clang/lib/Basic/Targets/NVPTX.cpp
+++ b/clang/lib/Basic/Targets/NVPTX.cpp
@@ -168,7 +168,7 @@ void NVPTXTargetInfo::getTargetDefines(const LangOptions &Opts,
MacroBuilder &Builder) const {
Builder.defineMacro("__PTX__");
Builder.defineMacro("__NVPTX__");
- if (Opts.CUDAIsDevice || Opts.OpenMPIsDevice || !HostTarget) {
+ if (Opts.CUDAIsDevice || Opts.OpenMPIsTargetDevice || !HostTarget) {
// Set __CUDA_ARCH__ for the GPU specified.
std::string CUDAArchCode = [this] {
switch (GPU) {
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index 8b503ef..7ea166d 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -5296,7 +5296,7 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
case Builtin::BIprintf:
if (getTarget().getTriple().isNVPTX() ||
getTarget().getTriple().isAMDGCN()) {
- if (getLangOpts().OpenMPIsDevice)
+ if (getLangOpts().OpenMPIsTargetDevice)
return EmitOpenMPDevicePrintfCallExpr(E);
if (getTarget().getTriple().isNVPTX())
return EmitNVPTXDevicePrintfCallExpr(E);
diff --git a/clang/lib/CodeGen/CGDecl.cpp b/clang/lib/CodeGen/CGDecl.cpp
index f19006f..f0953ee 100644
--- a/clang/lib/CodeGen/CGDecl.cpp
+++ b/clang/lib/CodeGen/CGDecl.cpp
@@ -1599,7 +1599,7 @@ CodeGenFunction::EmitAutoVarAlloca(const VarDecl &D) {
// OpenMP we have to use a call to __kmpc_alloc_shared(). The matching
// deallocation call to __kmpc_free_shared() is emitted later.
bool VarAllocated = false;
- if (getLangOpts().OpenMPIsDevice) {
+ if (getLangOpts().OpenMPIsTargetDevice) {
auto &RT = CGM.getOpenMPRuntime();
if (RT.isDelayedVariableLengthDecl(*this, &D)) {
// Emit call to __kmpc_alloc_shared() instead of the alloca.
diff --git a/clang/lib/CodeGen/CGExpr.cpp b/clang/lib/CodeGen/CGExpr.cpp
index 79b21bf..ae5168d 100644
--- a/clang/lib/CodeGen/CGExpr.cpp
+++ b/clang/lib/CodeGen/CGExpr.cpp
@@ -2588,7 +2588,7 @@ static LValue EmitGlobalVarDeclLValue(CodeGenFunction &CGF,
return CGF.CGM.getCXXABI().EmitThreadLocalVarDeclLValue(CGF, VD, T);
// Check if the variable is marked as declare target with link clause in
// device codegen.
- if (CGF.getLangOpts().OpenMPIsDevice) {
+ if (CGF.getLangOpts().OpenMPIsTargetDevice) {
Address Addr = emitDeclTargetVarDeclLValue(CGF, VD, T);
if (Addr.isValid())
return CGF.MakeAddrLValue(Addr, T, AlignmentSource::Decl);
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 4c32dc1..e5952bf 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -1054,10 +1054,10 @@ static FieldDecl *addFieldToRecordDecl(ASTContext &C, DeclContext *DC,
CGOpenMPRuntime::CGOpenMPRuntime(CodeGenModule &CGM)
: CGM(CGM), OMPBuilder(CGM.getModule()) {
KmpCriticalNameTy = llvm::ArrayType::get(CGM.Int32Ty, /*NumElements*/ 8);
- llvm::OpenMPIRBuilderConfig Config(CGM.getLangOpts().OpenMPIsDevice, false,
- hasRequiresUnifiedSharedMemory(),
+ llvm::OpenMPIRBuilderConfig Config(CGM.getLangOpts().OpenMPIsTargetDevice,
+ isGPU(), hasRequiresUnifiedSharedMemory(),
CGM.getLangOpts().OpenMPOffloadMandatory);
- OMPBuilder.initialize(CGM.getLangOpts().OpenMPIsDevice
+ OMPBuilder.initialize(CGM.getLangOpts().OpenMPIsTargetDevice
? CGM.getLangOpts().OMPHostIRFile
: StringRef{});
OMPBuilder.setConfig(Config);
@@ -1865,7 +1865,7 @@ bool CGOpenMPRuntime::emitDeclareTargetVarDefinition(const VarDecl *VD,
llvm::GlobalVariable *Addr,
bool PerformInit) {
if (CGM.getLangOpts().OMPTargetTriples.empty() &&
- !CGM.getLangOpts().OpenMPIsDevice)
+ !CGM.getLangOpts().OpenMPIsTargetDevice)
return false;
std::optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD);
@@ -1873,12 +1873,12 @@ bool CGOpenMPRuntime::emitDeclareTargetVarDefinition(const VarDecl *VD,
((*Res == OMPDeclareTargetDeclAttr::MT_To ||
*Res == OMPDeclareTargetDeclAttr::MT_Enter) &&
HasRequiresUnifiedSharedMemory))
- return CGM.getLangOpts().OpenMPIsDevice;
+ return CGM.getLangOpts().OpenMPIsTargetDevice;
VD = VD->getDefinition(CGM.getContext());
assert(VD && "Unknown VarDecl");
if (!DeclareTargetWithDefinition.insert(CGM.getMangledName(VD)).second)
- return CGM.getLangOpts().OpenMPIsDevice;
+ return CGM.getLangOpts().OpenMPIsTargetDevice;
QualType ASTTy = VD->getType();
SourceLocation Loc = VD->getCanonicalDecl()->getBeginLoc();
@@ -1895,7 +1895,7 @@ bool CGOpenMPRuntime::emitDeclareTargetVarDefinition(const VarDecl *VD,
if (CGM.getLangOpts().CPlusPlus && PerformInit) {
llvm::Constant *Ctor;
llvm::Constant *ID;
- if (CGM.getLangOpts().OpenMPIsDevice) {
+ if (CGM.getLangOpts().OpenMPIsTargetDevice) {
// Generate function that re-emits the declaration's initializer into
// the threadprivate copy of the variable VD
CodeGenFunction CtorCGF(CGM);
@@ -1943,7 +1943,7 @@ bool CGOpenMPRuntime::emitDeclareTargetVarDefinition(const VarDecl *VD,
if (VD->getType().isDestructedType() != QualType::DK_none) {
llvm::Constant *Dtor;
llvm::Constant *ID;
- if (CGM.getLangOpts().OpenMPIsDevice) {
+ if (CGM.getLangOpts().OpenMPIsTargetDevice) {
// Generate function that emits destructor call for the threadprivate
// copy of the variable VD
CodeGenFunction DtorCGF(CGM);
@@ -1988,7 +1988,7 @@ bool CGOpenMPRuntime::emitDeclareTargetVarDefinition(const VarDecl *VD,
DtorEntryInfo, Dtor, ID,
llvm::OffloadEntriesInfoManager::OMPTargetRegionEntryDtor);
}
- return CGM.getLangOpts().OpenMPIsDevice;
+ return CGM.getLangOpts().OpenMPIsTargetDevice;
}
Address CGOpenMPRuntime::getAddrOfArtificialThreadPrivate(CodeGenFunction &CGF,
@@ -2837,7 +2837,7 @@ void CGOpenMPRuntime::emitDistributeStaticInit(
llvm::Value *ThreadId = getThreadID(CGF, Loc);
llvm::FunctionCallee StaticInitFunction;
bool isGPUDistribute =
- CGM.getLangOpts().OpenMPIsDevice &&
+ CGM.getLangOpts().OpenMPIsTargetDevice &&
(CGM.getTriple().isAMDGCN() || CGM.getTriple().isNVPTX());
StaticInitFunction = createForStaticInitFunction(
Values.IVSize, Values.IVSigned, isGPUDistribute);
@@ -2862,7 +2862,8 @@ void CGOpenMPRuntime::emitForStaticFinish(CodeGenFunction &CGF,
: OMP_IDENT_WORK_SECTIONS),
getThreadID(CGF, Loc)};
auto DL = ApplyDebugLocation::CreateDefaultArtificial(CGF, Loc);
- if (isOpenMPDistributeDirective(DKind) && CGM.getLangOpts().OpenMPIsDevice &&
+ if (isOpenMPDistributeDirective(DKind) &&
+ CGM.getLangOpts().OpenMPIsTargetDevice &&
(CGM.getTriple().isAMDGCN() || CGM.getTriple().isNVPTX()))
CGF.EmitRuntimeCall(
OMPBuilder.getOrCreateRuntimeFunction(
@@ -6299,7 +6300,7 @@ const Expr *CGOpenMPRuntime::getNumTeamsExprForTargetDirective(
llvm::Value *CGOpenMPRuntime::emitNumTeamsForTargetDirective(
CodeGenFunction &CGF, const OMPExecutableDirective &D) {
- assert(!CGF.getLangOpts().OpenMPIsDevice &&
+ assert(!CGF.getLangOpts().OpenMPIsTargetDevice &&
"Clauses associated with the teams directive expected to be emitted "
"only for the host!");
CGBuilderTy &Bld = CGF.Builder;
@@ -6553,7 +6554,7 @@ const Expr *CGOpenMPRuntime::getNumThreadsExprForTargetDirective(
llvm::Value *CGOpenMPRuntime::emitNumThreadsForTargetDirective(
CodeGenFunction &CGF, const OMPExecutableDirective &D) {
- assert(!CGF.getLangOpts().OpenMPIsDevice &&
+ assert(!CGF.getLangOpts().OpenMPIsTargetDevice &&
"Clauses associated with the teams directive expected to be emitted "
"only for the host!");
OpenMPDirectiveKind DirectiveKind = D.getDirectiveKind();
@@ -9833,7 +9834,7 @@ void CGOpenMPRuntime::emitTargetCall(
if (!CGF.HaveInsertPoint())
return;
- const bool OffloadingMandatory = !CGM.getLangOpts().OpenMPIsDevice &&
+ const bool OffloadingMandatory = !CGM.getLangOpts().OpenMPIsTargetDevice &&
CGM.getLangOpts().OpenMPOffloadMandatory;
assert((OffloadingMandatory || OutlinedFn) && "Invalid outlined function!");
@@ -10061,10 +10062,10 @@ static bool isAssumedToBeNotEmitted(const ValueDecl *VD, bool IsDevice) {
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().OpenMPIsTargetDevice) {
if (const auto *FD = dyn_cast<FunctionDecl>(GD.getDecl()))
if (isAssumedToBeNotEmitted(cast<ValueDecl>(FD),
- CGM.getLangOpts().OpenMPIsDevice))
+ CGM.getLangOpts().OpenMPIsTargetDevice))
return true;
return false;
}
@@ -10075,7 +10076,7 @@ bool CGOpenMPRuntime::emitTargetFunctions(GlobalDecl GD) {
StringRef Name = CGM.getMangledName(GD);
scanForTargetRegionsFunctions(FD->getBody(), Name);
if (isAssumedToBeNotEmitted(cast<ValueDecl>(FD),
- CGM.getLangOpts().OpenMPIsDevice))
+ CGM.getLangOpts().OpenMPIsTargetDevice))
return true;
}
@@ -10086,10 +10087,10 @@ bool CGOpenMPRuntime::emitTargetFunctions(GlobalDecl GD) {
bool CGOpenMPRuntime::emitTargetGlobalVariable(GlobalDecl GD) {
if (isAssumedToBeNotEmitted(cast<ValueDecl>(GD.getDecl()),
- CGM.getLangOpts().OpenMPIsDevice))
+ CGM.getLangOpts().OpenMPIsTargetDevice))
return true;
- if (!CGM.getLangOpts().OpenMPIsDevice)
+ if (!CGM.getLangOpts().OpenMPIsTargetDevice)
return false;
// Check if there are Ctors/Dtors in this declaration and look for target
@@ -10126,13 +10127,13 @@ bool CGOpenMPRuntime::emitTargetGlobalVariable(GlobalDecl GD) {
void CGOpenMPRuntime::registerTargetGlobalVariable(const VarDecl *VD,
llvm::Constant *Addr) {
if (CGM.getLangOpts().OMPTargetTriples.empty() &&
- !CGM.getLangOpts().OpenMPIsDevice)
+ !CGM.getLangOpts().OpenMPIsTargetDevice)
return;
std::optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD);
if (!Res) {
- if (CGM.getLangOpts().OpenMPIsDevice) {
+ if (CGM.getLangOpts().OpenMPIsTargetDevice) {
// Register non-target variables being emitted in device code (debug info
// may cause this).
StringRef VarName = CGM.getMangledName(VD);
@@ -10260,19 +10261,19 @@ bool CGOpenMPRuntime::hasRequiresUnifiedSharedMemory() const {
CGOpenMPRuntime::DisableAutoDeclareTargetRAII::DisableAutoDeclareTargetRAII(
CodeGenModule &CGM)
: CGM(CGM) {
- if (CGM.getLangOpts().OpenMPIsDevice) {
+ if (CGM.getLangOpts().OpenMPIsTargetDevice) {
SavedShouldMarkAsGlobal = CGM.getOpenMPRuntime().ShouldMarkAsGlobal;
CGM.getOpenMPRuntime().ShouldMarkAsGlobal = false;
}
}
CGOpenMPRuntime::DisableAutoDeclareTargetRAII::~DisableAutoDeclareTargetRAII() {
- if (CGM.getLangOpts().OpenMPIsDevice)
+ if (CGM.getLangOpts().OpenMPIsTargetDevice)
CGM.getOpenMPRuntime().ShouldMarkAsGlobal = SavedShouldMarkAsGlobal;
}
bool CGOpenMPRuntime::markAsGlobalTarget(GlobalDecl GD) {
- if (!CGM.getLangOpts().OpenMPIsDevice || !ShouldMarkAsGlobal)
+ if (!CGM.getLangOpts().OpenMPIsTargetDevice || !ShouldMarkAsGlobal)
return true;
const auto *D = cast<FunctionDecl>(GD.getDecl());
@@ -10295,7 +10296,7 @@ llvm::Function *CGOpenMPRuntime::emitRequiresDirectiveRegFun() {
// If we don't have entries or if we are emitting code for the device, we
// don't need to do anything.
if (CGM.getLangOpts().OMPTargetTriples.empty() ||
- CGM.getLangOpts().OpenMPSimd || CGM.getLangOpts().OpenMPIsDevice ||
+ CGM.getLangOpts().OpenMPSimd || CGM.getLangOpts().OpenMPIsTargetDevice ||
(OMPBuilder.OffloadInfoManager.empty() &&
!HasEmittedDeclareTargetRegion && !HasEmittedTargetRegion))
return nullptr;
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.h b/clang/lib/CodeGen/CGOpenMPRuntime.h
index af801e3..ebdbf4c 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.h
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.h
@@ -688,7 +688,7 @@ public:
SizeEmitter);
/// Returns true if the current target is a GPU.
- virtual bool isTargetCodegen() const { return false; }
+ virtual bool isGPU() const { return false; }
/// Check if the variable length declaration is delayed:
virtual bool isDelayedVariableLengthDecl(CodeGenFunction &CGF,
diff --git a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
index ff71f57..e83514a3 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
@@ -873,12 +873,12 @@ void CGOpenMPRuntimeGPU::emitTargetOutlinedFunction(
CGOpenMPRuntimeGPU::CGOpenMPRuntimeGPU(CodeGenModule &CGM)
: CGOpenMPRuntime(CGM) {
- llvm::OpenMPIRBuilderConfig Config(CGM.getLangOpts().OpenMPIsDevice, true,
- hasRequiresUnifiedSharedMemory(),
+ llvm::OpenMPIRBuilderConfig Config(CGM.getLangOpts().OpenMPIsTargetDevice,
+ isGPU(), hasRequiresUnifiedSharedMemory(),
CGM.getLangOpts().OpenMPOffloadMandatory);
OMPBuilder.setConfig(Config);
- if (!CGM.getLangOpts().OpenMPIsDevice)
+ if (!CGM.getLangOpts().OpenMPIsTargetDevice)
llvm_unreachable("OpenMP can only handle device code.");
llvm::OpenMPIRBuilder &OMPBuilder = getOMPBuilder();
diff --git a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.h b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.h
index 59b071e..dddfe5a 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.h
+++ b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.h
@@ -119,7 +119,7 @@ public:
explicit CGOpenMPRuntimeGPU(CodeGenModule &CGM);
void clear() override;
- bool isTargetCodegen() const override { return true; };
+ bool isGPU() const override { return true; };
/// Declare generalized virtual functions which need to be defined
/// by all specializations of OpenMPGPURuntime Targets like AMDGCN
diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp
index d419896..74a0499 100644
--- a/clang/lib/CodeGen/CGStmtOpenMP.cpp
+++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp
@@ -802,7 +802,7 @@ bool CodeGenFunction::EmitOMPFirstprivateClause(const OMPExecutableDirective &D,
if (!HaveInsertPoint())
return false;
bool DeviceConstTarget =
- getLangOpts().OpenMPIsDevice &&
+ getLangOpts().OpenMPIsTargetDevice &&
isOpenMPTargetExecutionDirective(D.getDirectiveKind());
bool FirstprivateIsLastprivate = false;
llvm::DenseMap<const VarDecl *, OpenMPLastprivateModifier> Lastprivates;
@@ -6577,7 +6577,7 @@ static void emitCommonOMPTargetDirective(CodeGenFunction &CGF,
CodeGenModule &CGM = CGF.CGM;
// On device emit this construct as inlined code.
- if (CGM.getLangOpts().OpenMPIsDevice) {
+ if (CGM.getLangOpts().OpenMPIsTargetDevice) {
OMPLexicalScope Scope(CGF, S, OMPD_target);
CGM.getOpenMPRuntime().emitInlinedDirective(
CGF, OMPD_target, [&S](CodeGenFunction &CGF, PrePostActionTy &) {
diff --git a/clang/lib/CodeGen/CGVTables.cpp b/clang/lib/CodeGen/CGVTables.cpp
index fb0276a..409e294 100644
--- a/clang/lib/CodeGen/CGVTables.cpp
+++ b/clang/lib/CodeGen/CGVTables.cpp
@@ -775,7 +775,7 @@ void CodeGenVTables::addVTableComponent(ConstantArrayBuilder &builder,
// For NVPTX devices in OpenMP emit special functon as null pointers,
// otherwise linking ends up with unresolved references.
- if (CGM.getLangOpts().OpenMP && CGM.getLangOpts().OpenMPIsDevice &&
+ if (CGM.getLangOpts().OpenMP && CGM.getLangOpts().OpenMPIsTargetDevice &&
CGM.getTriple().isNVPTX())
return llvm::ConstantPointerNull::get(CGM.Int8PtrTy);
llvm::FunctionType *fnTy =
diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp
index 304646c..13b49eb 100644
--- a/clang/lib/CodeGen/CodeGenModule.cpp
+++ b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -470,7 +470,7 @@ void CodeGenModule::createOpenMPRuntime() {
case llvm::Triple::nvptx:
case llvm::Triple::nvptx64:
case llvm::Triple::amdgcn:
- assert(getLangOpts().OpenMPIsDevice &&
+ assert(getLangOpts().OpenMPIsTargetDevice &&
"OpenMP AMDGPU/NVPTX is only prepared to deal with device code.");
OpenMPRuntime.reset(new CGOpenMPRuntimeGPU(*this));
break;
@@ -1089,7 +1089,7 @@ void CodeGenModule::Release() {
// Indicate whether this Module was compiled with -fopenmp
if (getLangOpts().OpenMP && !getLangOpts().OpenMPSimd)
getModule().addModuleFlag(llvm::Module::Max, "openmp", LangOpts.OpenMP);
- if (getLangOpts().OpenMPIsDevice)
+ if (getLangOpts().OpenMPIsTargetDevice)
getModule().addModuleFlag(llvm::Module::Max, "openmp-device",
LangOpts.OpenMP);
@@ -4257,7 +4257,7 @@ llvm::Constant *CodeGenModule::GetOrCreateLLVMFunction(
// the iFunc instead. Name Mangling will handle the rest of the changes.
if (const FunctionDecl *FD = cast_or_null<FunctionDecl>(D)) {
// For the device mark the function as one that should be emitted.
- if (getLangOpts().OpenMPIsDevice && OpenMPRuntime &&
+ if (getLangOpts().OpenMPIsTargetDevice && OpenMPRuntime &&
!OpenMPRuntime->markAsGlobalTarget(GD) && FD->isDefined() &&
!DontDefer && !IsForDefinition) {
if (const FunctionDecl *FDDef = FD->getDefinition()) {
@@ -5081,7 +5081,7 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D,
// If this is OpenMP device, check if it is legal to emit this global
// normally.
- if (LangOpts.OpenMPIsDevice && OpenMPRuntime &&
+ if (LangOpts.OpenMPIsTargetDevice && OpenMPRuntime &&
OpenMPRuntime->emitTargetGlobalVariable(D))
return;
@@ -6717,7 +6717,7 @@ void CodeGenModule::EmitTopLevelDecl(Decl *D) {
if (LangOpts.CUDA && LangOpts.CUDAIsDevice)
break;
// File-scope asm is ignored during device-side OpenMP compilation.
- if (LangOpts.OpenMPIsDevice)
+ if (LangOpts.OpenMPIsTargetDevice)
break;
// File-scope asm is ignored during device-side SYCL compilation.
if (LangOpts.SYCLIsDevice)
@@ -7148,7 +7148,7 @@ llvm::Constant *CodeGenModule::GetAddrOfRTTIDescriptor(QualType Ty,
// FIXME: should we even be calling this method if RTTI is disabled
// and it's not for EH?
if ((!ForEH && !getLangOpts().RTTI) || getLangOpts().CUDAIsDevice ||
- (getLangOpts().OpenMP && getLangOpts().OpenMPIsDevice &&
+ (getLangOpts().OpenMP && getLangOpts().OpenMPIsTargetDevice &&
getTriple().isNVPTX()))
return llvm::Constant::getNullValue(Int8PtrTy);
diff --git a/clang/lib/CodeGen/Targets/NVPTX.cpp b/clang/lib/CodeGen/Targets/NVPTX.cpp
index 75fda51..1ca0192 100644
--- a/clang/lib/CodeGen/Targets/NVPTX.cpp
+++ b/clang/lib/CodeGen/Targets/NVPTX.cpp
@@ -149,7 +149,8 @@ ABIArgInfo NVPTXABIInfo::classifyReturnType(QualType RetTy) const {
return ABIArgInfo::getIgnore();
if (getContext().getLangOpts().OpenMP &&
- getContext().getLangOpts().OpenMPIsDevice && isUnsupportedType(RetTy))
+ getContext().getLangOpts().OpenMPIsTargetDevice &&
+ isUnsupportedType(RetTy))
return coerceToIntArrayWithLimit(RetTy, 64);
// note: this is different from default ABI
diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp
index e3d7b05..efdb788 100644
--- a/clang/lib/Driver/ToolChains/Clang.cpp
+++ b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -7239,11 +7239,11 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
// OpenMP offloading device jobs take the argument -fopenmp-host-ir-file-path
// to specify the result of the compile phase on the host, so the meaningful
- // device declarations can be identified. Also, -fopenmp-is-device is passed
- // along to tell the frontend that it is generating code for a device, so that
- // only the relevant declarations are emitted.
+ // device declarations can be identified. Also, -fopenmp-is-target-device is
+ // passed along to tell the frontend that it is generating code for a device,
+ // so that only the relevant declarations are emitted.
if (IsOpenMPDevice) {
- CmdArgs.push_back("-fopenmp-is-device");
+ CmdArgs.push_back("-fopenmp-is-target-device");
if (OpenMPDeviceInput) {
CmdArgs.push_back("-fopenmp-host-ir-file-path");
CmdArgs.push_back(Args.MakeArgString(OpenMPDeviceInput->getFilename()));
diff --git a/clang/lib/Driver/ToolChains/Flang.cpp b/clang/lib/Driver/ToolChains/Flang.cpp
index 75b9524..bdc75f5 100644
--- a/clang/lib/Driver/ToolChains/Flang.cpp
+++ b/clang/lib/Driver/ToolChains/Flang.cpp
@@ -231,10 +231,9 @@ void Flang::addOffloadOptions(Compilation &C, const InputInfoList &Inputs,
}
if (IsOpenMPDevice) {
- // -fopenmp-is-device is passed along to tell the frontend that it is
- // generating code for a device, so that only the relevant code is
- // emitted.
- CmdArgs.push_back("-fopenmp-is-device");
+ // -fopenmp-is-target-device is passed along to tell the frontend that it is
+ // generating code for a device, so that only the relevant code is emitted.
+ CmdArgs.push_back("-fopenmp-is-target-device");
// When in OpenMP offloading mode, enable debugging on the device.
Args.AddAllArgs(CmdArgs, options::OPT_fopenmp_target_debug_EQ);
diff --git a/clang/lib/Frontend/CompilerInstance.cpp b/clang/lib/Frontend/CompilerInstance.cpp
index fc1295a..a87e91e 100644
--- a/clang/lib/Frontend/CompilerInstance.cpp
+++ b/clang/lib/Frontend/CompilerInstance.cpp
@@ -113,7 +113,7 @@ bool CompilerInstance::createTarget() {
// Check whether AuxTarget exists, if not, then create TargetInfo for the
// other side of CUDA/OpenMP/SYCL compilation.
if (!getAuxTarget() &&
- (getLangOpts().CUDA || getLangOpts().OpenMPIsDevice ||
+ (getLangOpts().CUDA || getLangOpts().OpenMPIsTargetDevice ||
getLangOpts().SYCLIsDevice) &&
!getFrontendOpts().AuxTriple.empty()) {
auto TO = std::make_shared<TargetOptions>();
diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp
index 50440e6..5360bbe 100644
--- a/clang/lib/Frontend/CompilerInvocation.cpp
+++ b/clang/lib/Frontend/CompilerInvocation.cpp
@@ -3401,8 +3401,8 @@ void CompilerInvocation::GenerateLangArgs(const LangOptions &Opts,
if (!Opts.OpenMPUseTLS)
GenerateArg(Args, OPT_fnoopenmp_use_tls, SA);
- if (Opts.OpenMPIsDevice)
- GenerateArg(Args, OPT_fopenmp_is_device, SA);
+ if (Opts.OpenMPIsTargetDevice)
+ GenerateArg(Args, OPT_fopenmp_is_target_device, SA);
if (Opts.OpenMPIRBuilder)
GenerateArg(Args, OPT_fopenmp_enable_irbuilder, SA);
@@ -3787,14 +3787,15 @@ bool CompilerInvocation::ParseLangArgs(LangOptions &Opts, ArgList &Args,
Opts.OpenMPSimd = !Opts.OpenMP && IsSimdSpecified;
Opts.OpenMPUseTLS =
Opts.OpenMP && !Args.hasArg(options::OPT_fnoopenmp_use_tls);
- Opts.OpenMPIsDevice =
- Opts.OpenMP && Args.hasArg(options::OPT_fopenmp_is_device);
+ Opts.OpenMPIsTargetDevice =
+ Opts.OpenMP && Args.hasArg(options::OPT_fopenmp_is_target_device);
Opts.OpenMPIRBuilder =
Opts.OpenMP && Args.hasArg(options::OPT_fopenmp_enable_irbuilder);
bool IsTargetSpecified =
- Opts.OpenMPIsDevice || Args.hasArg(options::OPT_fopenmp_targets_EQ);
+ Opts.OpenMPIsTargetDevice || Args.hasArg(options::OPT_fopenmp_targets_EQ);
- Opts.ConvergentFunctions = Opts.ConvergentFunctions || Opts.OpenMPIsDevice;
+ Opts.ConvergentFunctions =
+ Opts.ConvergentFunctions || Opts.OpenMPIsTargetDevice;
if (Opts.OpenMP || Opts.OpenMPSimd) {
if (int Version = getLastArgIntValue(
@@ -3803,7 +3804,7 @@ bool CompilerInvocation::ParseLangArgs(LangOptions &Opts, ArgList &Args,
Opts.OpenMP = Version;
// Provide diagnostic when a given target is not expected to be an OpenMP
// device or host.
- if (!Opts.OpenMPIsDevice) {
+ if (!Opts.OpenMPIsTargetDevice) {
switch (T.getArch()) {
default:
break;
@@ -3818,13 +3819,13 @@ bool CompilerInvocation::ParseLangArgs(LangOptions &Opts, ArgList &Args,
// Set the flag to prevent the implementation from emitting device exception
// handling code for those requiring so.
- if ((Opts.OpenMPIsDevice && (T.isNVPTX() || T.isAMDGCN())) ||
+ if ((Opts.OpenMPIsTargetDevice && (T.isNVPTX() || T.isAMDGCN())) ||
Opts.OpenCLCPlusPlus) {
Opts.Exceptions = 0;
Opts.CXXExceptions = 0;
}
- if (Opts.OpenMPIsDevice && T.isNVPTX()) {
+ if (Opts.OpenMPIsTargetDevice && T.isNVPTX()) {
Opts.OpenMPCUDANumSMs =
getLastArgIntValue(Args, options::OPT_fopenmp_cuda_number_of_sm_EQ,
Opts.OpenMPCUDANumSMs, Diags);
@@ -3838,15 +3839,15 @@ bool CompilerInvocation::ParseLangArgs(LangOptions &Opts, ArgList &Args,
// Set the value of the debugging flag used in the new offloading device RTL.
// Set either by a specific value or to a default if not specified.
- if (Opts.OpenMPIsDevice && (Args.hasArg(OPT_fopenmp_target_debug) ||
- Args.hasArg(OPT_fopenmp_target_debug_EQ))) {
+ if (Opts.OpenMPIsTargetDevice && (Args.hasArg(OPT_fopenmp_target_debug) ||
+ Args.hasArg(OPT_fopenmp_target_debug_EQ))) {
Opts.OpenMPTargetDebug = getLastArgIntValue(
Args, OPT_fopenmp_target_debug_EQ, Opts.OpenMPTargetDebug, Diags);
if (!Opts.OpenMPTargetDebug && Args.hasArg(OPT_fopenmp_target_debug))
Opts.OpenMPTargetDebug = 1;
}
- if (Opts.OpenMPIsDevice) {
+ if (Opts.OpenMPIsTargetDevice) {
if (Args.hasArg(OPT_fopenmp_assume_teams_oversubscription))
Opts.OpenMPTeamSubscription = true;
if (Args.hasArg(OPT_fopenmp_assume_threads_oversubscription))
@@ -3894,7 +3895,8 @@ bool CompilerInvocation::ParseLangArgs(LangOptions &Opts, ArgList &Args,
}
// Set CUDA mode for OpenMP target NVPTX/AMDGCN if specified in options
- Opts.OpenMPCUDAMode = Opts.OpenMPIsDevice && (T.isNVPTX() || T.isAMDGCN()) &&
+ Opts.OpenMPCUDAMode = Opts.OpenMPIsTargetDevice &&
+ (T.isNVPTX() || T.isAMDGCN()) &&
Args.hasArg(options::OPT_fopenmp_cuda_mode);
// FIXME: Eliminate this dependency.
@@ -4432,7 +4434,7 @@ bool CompilerInvocation::CreateFromArgsImpl(
}
// Set the triple of the host for OpenMP device compile.
- if (LangOpts.OpenMPIsDevice)
+ if (LangOpts.OpenMPIsTargetDevice)
Res.getTargetOpts().HostTriple = Res.getFrontendOpts().AuxTriple;
ParseCodeGenArgs(Res.getCodeGenOpts(), Args, DashX, Diags, T,
diff --git a/clang/lib/Frontend/InitPreprocessor.cpp b/clang/lib/Frontend/InitPreprocessor.cpp
index 431b9ea..9a83bec 100644
--- a/clang/lib/Frontend/InitPreprocessor.cpp
+++ b/clang/lib/Frontend/InitPreprocessor.cpp
@@ -1343,7 +1343,8 @@ void clang::InitializePreprocessor(
if (InitOpts.UsePredefines) {
// FIXME: This will create multiple definitions for most of the predefined
// macros. This is not the right way to handle this.
- if ((LangOpts.CUDA || LangOpts.OpenMPIsDevice || LangOpts.SYCLIsDevice) &&
+ if ((LangOpts.CUDA || LangOpts.OpenMPIsTargetDevice ||
+ LangOpts.SYCLIsDevice) &&
PP.getAuxTargetInfo())
InitializePredefinedMacros(*PP.getAuxTargetInfo(), LangOpts, FEOpts,
PP.getPreprocessorOpts(), Builder);
diff --git a/clang/lib/Lex/LiteralSupport.cpp b/clang/lib/Lex/LiteralSupport.cpp
index 736d08d..2ca9946 100644
--- a/clang/lib/Lex/LiteralSupport.cpp
+++ b/clang/lib/Lex/LiteralSupport.cpp
@@ -975,7 +975,7 @@ NumericLiteralParser::NumericLiteralParser(StringRef TokSpelling,
// ToDo: more precise check for CUDA.
// TODO: AMDGPU might also support it in the future.
if ((Target.hasFloat16Type() || LangOpts.CUDA ||
- (LangOpts.OpenMPIsDevice && Target.getTriple().isNVPTX())) &&
+ (LangOpts.OpenMPIsTargetDevice && Target.getTriple().isNVPTX())) &&
s + 2 < ThisTokEnd && s[1] == '1' && s[2] == '6') {
s += 2; // success, eat up 2 characters.
isFloat16 = true;
diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp
index 92e976b..46ae6fb 100644
--- a/clang/lib/Sema/Sema.cpp
+++ b/clang/lib/Sema/Sema.cpp
@@ -1880,8 +1880,9 @@ Sema::SemaDiagnosticBuilder
Sema::targetDiag(SourceLocation Loc, unsigned DiagID, const FunctionDecl *FD) {
FD = FD ? FD : getCurFunctionDecl();
if (LangOpts.OpenMP)
- return LangOpts.OpenMPIsDevice ? diagIfOpenMPDeviceCode(Loc, DiagID, FD)
- : diagIfOpenMPHostCode(Loc, DiagID, FD);
+ return LangOpts.OpenMPIsTargetDevice
+ ? diagIfOpenMPDeviceCode(Loc, DiagID, FD)
+ : diagIfOpenMPHostCode(Loc, DiagID, FD);
if (getLangOpts().CUDA)
return getLangOpts().CUDAIsDevice ? CUDADiagIfDeviceCode(Loc, DiagID)
: CUDADiagIfHostCode(Loc, DiagID);
@@ -2006,7 +2007,8 @@ void Sema::checkTypeSupport(QualType Ty, SourceLocation Loc, ValueDecl *D) {
};
auto CheckType = [&](QualType Ty, bool IsRetTy = false) {
- if (LangOpts.SYCLIsDevice || (LangOpts.OpenMP && LangOpts.OpenMPIsDevice) ||
+ if (LangOpts.SYCLIsDevice ||
+ (LangOpts.OpenMP && LangOpts.OpenMPIsTargetDevice) ||
LangOpts.CUDAIsDevice)
CheckDeviceType(Ty);
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 7a4e154..0cbf608 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -7763,7 +7763,7 @@ NamedDecl *Sema::ActOnVariableDeclarator(
diag::err_thread_non_global)
<< DeclSpec::getSpecifierName(TSCS);
else if (!Context.getTargetInfo().isTLSSupported()) {
- if (getLangOpts().CUDA || getLangOpts().OpenMPIsDevice ||
+ if (getLangOpts().CUDA || getLangOpts().OpenMPIsTargetDevice ||
getLangOpts().SYCLIsDevice) {
// Postpone error emission until we've collected attributes required to
// figure out whether it's a host or device variable and whether the
@@ -7892,17 +7892,18 @@ NamedDecl *Sema::ActOnVariableDeclarator(
if (const auto *TT = R->getAs<TypedefType>())
copyAttrFromTypedefToDecl<AllocSizeAttr>(*this, NewVD, TT);
- if (getLangOpts().CUDA || getLangOpts().OpenMPIsDevice ||
+ if (getLangOpts().CUDA || getLangOpts().OpenMPIsTargetDevice ||
getLangOpts().SYCLIsDevice) {
if (EmitTLSUnsupportedError &&
((getLangOpts().CUDA && DeclAttrsMatchCUDAMode(getLangOpts(), NewVD)) ||
- (getLangOpts().OpenMPIsDevice &&
+ (getLangOpts().OpenMPIsTargetDevice &&
OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(NewVD))))
Diag(D.getDeclSpec().getThreadStorageClassSpecLoc(),
diag::err_thread_unsupported);
if (EmitTLSUnsupportedError &&
- (LangOpts.SYCLIsDevice || (LangOpts.OpenMP && LangOpts.OpenMPIsDevice)))
+ (LangOpts.SYCLIsDevice ||
+ (LangOpts.OpenMP && LangOpts.OpenMPIsTargetDevice)))
targetDiag(D.getIdentifierLoc(), diag::err_thread_unsupported);
// CUDA B.2.5: "__shared__ and __constant__ variables have implied static
// storage [duration]."
@@ -13595,7 +13596,7 @@ void Sema::AddInitializerToDecl(Decl *RealDecl, Expr *Init, bool DirectInit) {
}
if (LangOpts.OpenMP &&
- (LangOpts.OpenMPIsDevice || !LangOpts.OMPTargetTriples.empty()) &&
+ (LangOpts.OpenMPIsTargetDevice || !LangOpts.OMPTargetTriples.empty()) &&
VDecl->isFileVarDecl())
DeclsToCheckForDeferredDiags.insert(VDecl);
CheckCompleteVariableDeclaration(VDecl);
@@ -15961,7 +15962,7 @@ Decl *Sema::ActOnFinishFunctionBody(Decl *dcl, Stmt *Body,
DiscardCleanupsInEvaluationContext();
}
- if (FD && ((LangOpts.OpenMP && (LangOpts.OpenMPIsDevice ||
+ if (FD && ((LangOpts.OpenMP && (LangOpts.OpenMPIsTargetDevice ||
!LangOpts.OMPTargetTriples.empty())) ||
LangOpts.CUDA || LangOpts.SYCLIsDevice)) {
auto ES = getEmissionStatus(FD);
@@ -20039,7 +20040,7 @@ Sema::FunctionEmissionStatus Sema::getEmissionStatus(const FunctionDecl *FD,
getASTContext().GetGVALinkageForFunction(Def));
};
- if (LangOpts.OpenMPIsDevice) {
+ if (LangOpts.OpenMPIsTargetDevice) {
// In OpenMP device mode we will not emit host only functions, or functions
// we don't need due to their linkage.
std::optional<OMPDeclareTargetDeclAttr::DevTypeTy> DevTy =
diff --git a/clang/lib/Sema/SemaDeclCXX.cpp b/clang/lib/Sema/SemaDeclCXX.cpp
index 507cfc4..5c68642 100644
--- a/clang/lib/Sema/SemaDeclCXX.cpp
+++ b/clang/lib/Sema/SemaDeclCXX.cpp
@@ -18065,7 +18065,7 @@ void Sema::MarkVTableUsed(SourceLocation Loc, CXXRecordDecl *Class,
return;
// Do not mark as used if compiling for the device outside of the target
// region.
- if (TUKind != TU_Prefix && LangOpts.OpenMP && LangOpts.OpenMPIsDevice &&
+ if (TUKind != TU_Prefix && LangOpts.OpenMP && LangOpts.OpenMPIsTargetDevice &&
!isInOpenMPDeclareTargetContext() &&
!isInOpenMPTargetExecutionDirective()) {
if (!DefinitionRequired)
diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp
index fa1f5e3..56e9c4c 100644
--- a/clang/lib/Sema/SemaExpr.cpp
+++ b/clang/lib/Sema/SemaExpr.cpp
@@ -387,7 +387,8 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef<SourceLocation> Locs,
if (auto *VD = dyn_cast<ValueDecl>(D))
checkTypeSupport(VD->getType(), Loc, VD);
- if (LangOpts.SYCLIsDevice || (LangOpts.OpenMP && LangOpts.OpenMPIsDevice)) {
+ if (LangOpts.SYCLIsDevice ||
+ (LangOpts.OpenMP && LangOpts.OpenMPIsTargetDevice)) {
if (!Context.getTargetInfo().isTLSSupported())
if (const auto *VD = dyn_cast<VarDecl>(D))
if (VD->getTLSKind() != VarDecl::TLS_None)
@@ -17163,7 +17164,7 @@ ExprResult Sema::BuildVAArgExpr(SourceLocation BuiltinLoc,
}
// NVPTX does not support va_arg expression.
- if (getLangOpts().OpenMP && getLangOpts().OpenMPIsDevice &&
+ if (getLangOpts().OpenMP && getLangOpts().OpenMPIsTargetDevice &&
Context.getTargetInfo().getTriple().isNVPTX())
targetDiag(E->getBeginLoc(), diag::err_va_arg_in_device);
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index cc52306d..970dd7b 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -2010,7 +2010,7 @@ void Sema::popOpenMPFunctionRegion(const FunctionScopeInfo *OldFSI) {
}
static bool isOpenMPDeviceDelayedContext(Sema &S) {
- assert(S.LangOpts.OpenMP && S.LangOpts.OpenMPIsDevice &&
+ assert(S.LangOpts.OpenMP && S.LangOpts.OpenMPIsTargetDevice &&
"Expected OpenMP device compilation.");
return !S.isInOpenMPTargetExecutionDirective();
}
@@ -2027,7 +2027,7 @@ enum class FunctionEmissionStatus {
Sema::SemaDiagnosticBuilder
Sema::diagIfOpenMPDeviceCode(SourceLocation Loc, unsigned DiagID,
const FunctionDecl *FD) {
- assert(LangOpts.OpenMP && LangOpts.OpenMPIsDevice &&
+ assert(LangOpts.OpenMP && LangOpts.OpenMPIsTargetDevice &&
"Expected OpenMP device compilation.");
SemaDiagnosticBuilder::Kind Kind = SemaDiagnosticBuilder::K_Nop;
@@ -2065,7 +2065,7 @@ Sema::diagIfOpenMPDeviceCode(SourceLocation Loc, unsigned DiagID,
Sema::SemaDiagnosticBuilder Sema::diagIfOpenMPHostCode(SourceLocation Loc,
unsigned DiagID,
const FunctionDecl *FD) {
- assert(LangOpts.OpenMP && !LangOpts.OpenMPIsDevice &&
+ assert(LangOpts.OpenMP && !LangOpts.OpenMPIsTargetDevice &&
"Expected OpenMP host compilation.");
SemaDiagnosticBuilder::Kind Kind = SemaDiagnosticBuilder::K_Nop;
@@ -2702,16 +2702,16 @@ void Sema::finalizeOpenMPDelayedAnalysis(const FunctionDecl *Caller,
std::optional<OMPDeclareTargetDeclAttr::DevTypeTy> DevTy =
OMPDeclareTargetDeclAttr::getDeviceType(Caller->getMostRecentDecl());
// Ignore host functions during device analyzis.
- if (LangOpts.OpenMPIsDevice &&
+ if (LangOpts.OpenMPIsTargetDevice &&
(!DevTy || *DevTy == OMPDeclareTargetDeclAttr::DT_Host))
return;
// Ignore nohost functions during host analyzis.
- if (!LangOpts.OpenMPIsDevice && DevTy &&
+ if (!LangOpts.OpenMPIsTargetDevice && DevTy &&
*DevTy == OMPDeclareTargetDeclAttr::DT_NoHost)
return;
const FunctionDecl *FD = Callee->getMostRecentDecl();
DevTy = OMPDeclareTargetDeclAttr::getDeviceType(FD);
- if (LangOpts.OpenMPIsDevice && DevTy &&
+ if (LangOpts.OpenMPIsTargetDevice && DevTy &&
*DevTy == OMPDeclareTargetDeclAttr::DT_Host) {
// Diagnose host function called during device codegen.
StringRef HostDevTy =
@@ -2722,8 +2722,8 @@ void Sema::finalizeOpenMPDelayedAnalysis(const FunctionDecl *Caller,
<< HostDevTy;
return;
}
- if (!LangOpts.OpenMPIsDevice && !LangOpts.OpenMPOffloadMandatory && DevTy &&
- *DevTy == OMPDeclareTargetDeclAttr::DT_NoHost) {
+ if (!LangOpts.OpenMPIsTargetDevice && !LangOpts.OpenMPOffloadMandatory &&
+ DevTy && *DevTy == OMPDeclareTargetDeclAttr::DT_NoHost) {
// In OpenMP 5.2 or later, if the function has a host variant then allow
// that to be called instead
auto &&HasHostAttr = [](const FunctionDecl *Callee) {
@@ -3386,7 +3386,7 @@ Sema::ActOnOpenMPAllocateDirective(SourceLocation Loc, ArrayRef<Expr *> VarList,
// allocate directives that appear in a target region must specify an
// allocator clause unless a requires directive with the dynamic_allocators
// clause is present in the same compilation unit.
- if (LangOpts.OpenMPIsDevice &&
+ if (LangOpts.OpenMPIsTargetDevice &&
!DSAStack->hasRequiresDeclWithClause<OMPDynamicAllocatorsClause>())
targetDiag(Loc, diag::err_expected_allocator_clause);
} else {
@@ -23659,7 +23659,7 @@ OMPClause *Sema::ActOnOpenMPAllocateClause(
// 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 &&
+ if (LangOpts.OpenMPIsTargetDevice &&
!DSAStack->hasRequiresDeclWithClause<OMPDynamicAllocatorsClause>())
targetDiag(StartLoc, diag::err_expected_allocator_expression);
}
diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp
index f5384e4..0aa691d 100644
--- a/clang/lib/Sema/SemaType.cpp
+++ b/clang/lib/Sema/SemaType.cpp
@@ -1502,7 +1502,7 @@ static QualType ConvertDeclSpecToType(TypeProcessingState &state) {
case DeclSpec::TST_int128:
if (!S.Context.getTargetInfo().hasInt128Type() &&
!(S.getLangOpts().SYCLIsDevice || S.getLangOpts().CUDAIsDevice ||
- (S.getLangOpts().OpenMP && S.getLangOpts().OpenMPIsDevice)))
+ (S.getLangOpts().OpenMP && S.getLangOpts().OpenMPIsTargetDevice)))
S.Diag(DS.getTypeSpecTypeLoc(), diag::err_type_unsupported)
<< "__int128";
if (DS.getTypeSpecSign() == TypeSpecifierSign::Unsigned)
@@ -1515,7 +1515,7 @@ static QualType ConvertDeclSpecToType(TypeProcessingState &state) {
// do not diagnose _Float16 usage to avoid false alarm.
// ToDo: more precise diagnostics for CUDA.
if (!S.Context.getTargetInfo().hasFloat16Type() && !S.getLangOpts().CUDA &&
- !(S.getLangOpts().OpenMP && S.getLangOpts().OpenMPIsDevice))
+ !(S.getLangOpts().OpenMP && S.getLangOpts().OpenMPIsTargetDevice))
S.Diag(DS.getTypeSpecTypeLoc(), diag::err_type_unsupported)
<< "_Float16";
Result = Context.Float16Ty;
@@ -1523,7 +1523,7 @@ static QualType ConvertDeclSpecToType(TypeProcessingState &state) {
case DeclSpec::TST_half: Result = Context.HalfTy; break;
case DeclSpec::TST_BFloat16:
if (!S.Context.getTargetInfo().hasBFloat16Type() &&
- !(S.getLangOpts().OpenMP && S.getLangOpts().OpenMPIsDevice) &&
+ !(S.getLangOpts().OpenMP && S.getLangOpts().OpenMPIsTargetDevice) &&
!S.getLangOpts().SYCLIsDevice)
S.Diag(DS.getTypeSpecTypeLoc(), diag::err_type_unsupported) << "__bf16";
Result = Context.BFloat16Ty;
@@ -1548,7 +1548,7 @@ static QualType ConvertDeclSpecToType(TypeProcessingState &state) {
case DeclSpec::TST_float128:
if (!S.Context.getTargetInfo().hasFloat128Type() &&
!S.getLangOpts().SYCLIsDevice &&
- !(S.getLangOpts().OpenMP && S.getLangOpts().OpenMPIsDevice))
+ !(S.getLangOpts().OpenMP && S.getLangOpts().OpenMPIsTargetDevice))
S.Diag(DS.getTypeSpecTypeLoc(), diag::err_type_unsupported)
<< "__float128";
Result = Context.Float128Ty;
@@ -1556,7 +1556,7 @@ static QualType ConvertDeclSpecToType(TypeProcessingState &state) {
case DeclSpec::TST_ibm128:
if (!S.Context.getTargetInfo().hasIbm128Type() &&
!S.getLangOpts().SYCLIsDevice &&
- !(S.getLangOpts().OpenMP && S.getLangOpts().OpenMPIsDevice))
+ !(S.getLangOpts().OpenMP && S.getLangOpts().OpenMPIsTargetDevice))
S.Diag(DS.getTypeSpecTypeLoc(), diag::err_type_unsupported) << "__ibm128";
Result = Context.Ibm128Ty;
break;
generated by cgit v1.1 at 2025-08-03 19:23:38 +0000