summaryrefslogtreecommitdiffstats
path: root/clang/lib
diff options
context:
space:
mode:
Diffstat (limited to 'clang/lib')
-rw-r--r--clang/lib/Parse/ParsePragma.cpp44
-rw-r--r--clang/lib/Sema/SemaCUDA.cpp27
-rw-r--r--clang/lib/Serialization/ASTReader.cpp9
-rw-r--r--clang/lib/Serialization/ASTWriter.cpp9
4 files changed, 89 insertions, 0 deletions
diff --git a/clang/lib/Parse/ParsePragma.cpp b/clang/lib/Parse/ParsePragma.cpp
index 7ae03af2b16..d6539c9610b 100644
--- a/clang/lib/Parse/ParsePragma.cpp
+++ b/clang/lib/Parse/ParsePragma.cpp
@@ -167,6 +167,16 @@ struct PragmaMSIntrinsicHandler : public PragmaHandler {
Token &FirstToken) override;
};
+struct PragmaForceCUDAHostDeviceHandler : public PragmaHandler {
+ PragmaForceCUDAHostDeviceHandler(Sema &Actions)
+ : PragmaHandler("force_cuda_host_device"), Actions(Actions) {}
+ void HandlePragma(Preprocessor &PP, PragmaIntroducerKind Introducer,
+ Token &FirstToken) override;
+
+private:
+ Sema &Actions;
+};
+
} // end namespace
void Parser::initializePragmaHandlers() {
@@ -239,6 +249,12 @@ void Parser::initializePragmaHandlers() {
PP.AddPragmaHandler(MSIntrinsic.get());
}
+ if (getLangOpts().CUDA) {
+ CUDAForceHostDeviceHandler.reset(
+ new PragmaForceCUDAHostDeviceHandler(Actions));
+ PP.AddPragmaHandler("clang", CUDAForceHostDeviceHandler.get());
+ }
+
OptimizeHandler.reset(new PragmaOptimizeHandler(Actions));
PP.AddPragmaHandler("clang", OptimizeHandler.get());
@@ -309,6 +325,11 @@ void Parser::resetPragmaHandlers() {
MSIntrinsic.reset();
}
+ if (getLangOpts().CUDA) {
+ PP.RemovePragmaHandler("clang", CUDAForceHostDeviceHandler.get());
+ CUDAForceHostDeviceHandler.reset();
+ }
+
PP.RemovePragmaHandler("STDC", FPContractHandler.get());
FPContractHandler.reset();
@@ -2187,3 +2208,26 @@ void PragmaMSIntrinsicHandler::HandlePragma(Preprocessor &PP,
PP.Diag(Tok.getLocation(), diag::warn_pragma_extra_tokens_at_eol)
<< "intrinsic";
}
+void PragmaForceCUDAHostDeviceHandler::HandlePragma(
+ Preprocessor &PP, PragmaIntroducerKind Introducer, Token &Tok) {
+ Token FirstTok = Tok;
+
+ PP.Lex(Tok);
+ IdentifierInfo *Info = Tok.getIdentifierInfo();
+ if (!Info || (!Info->isStr("begin") && !Info->isStr("end"))) {
+ PP.Diag(FirstTok.getLocation(),
+ diag::warn_pragma_force_cuda_host_device_bad_arg);
+ return;
+ }
+
+ if (Info->isStr("begin"))
+ Actions.PushForceCUDAHostDevice();
+ else if (!Actions.PopForceCUDAHostDevice())
+ PP.Diag(FirstTok.getLocation(),
+ diag::err_pragma_cannot_end_force_cuda_host_device);
+
+ PP.Lex(Tok);
+ if (!Tok.is(tok::eod))
+ PP.Diag(FirstTok.getLocation(),
+ diag::warn_pragma_force_cuda_host_device_bad_arg);
+}
diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index cb7019242f1..d6c0606674e 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -23,6 +23,19 @@
#include "llvm/ADT/SmallVector.h"
using namespace clang;
+void Sema::PushForceCUDAHostDevice() {
+ assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
+ ForceCUDAHostDeviceDepth++;
+}
+
+bool Sema::PopForceCUDAHostDevice() {
+ assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
+ if (ForceCUDAHostDeviceDepth == 0)
+ return false;
+ ForceCUDAHostDeviceDepth--;
+ return true;
+}
+
ExprResult Sema::ActOnCUDAExecConfigExpr(Scope *S, SourceLocation LLLLoc,
MultiExprArg ExecConfig,
SourceLocation GGGLoc) {
@@ -441,9 +454,23 @@ bool Sema::isEmptyCudaDestructor(SourceLocation Loc, CXXDestructorDecl *DD) {
// * a __device__ function with this signature was already declared, in which
// case in which case we output an error, unless the __device__ decl is in a
// system header, in which case we leave the constexpr function unattributed.
+//
+// In addition, all function decls are treated as __host__ __device__ when
+// ForceCUDAHostDeviceDepth > 0 (corresponding to code within a
+// #pragma clang force_cuda_host_device_begin/end
+// pair).
void Sema::maybeAddCUDAHostDeviceAttrs(Scope *S, FunctionDecl *NewD,
const LookupResult &Previous) {
assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
+
+ if (ForceCUDAHostDeviceDepth > 0) {
+ if (!NewD->hasAttr<CUDAHostAttr>())
+ NewD->addAttr(CUDAHostAttr::CreateImplicit(Context));
+ if (!NewD->hasAttr<CUDADeviceAttr>())
+ NewD->addAttr(CUDADeviceAttr::CreateImplicit(Context));
+ return;
+ }
+
if (!getLangOpts().CUDAHostDeviceConstexpr || !NewD->isConstexpr() ||
NewD->isVariadic() || NewD->hasAttr<CUDAHostAttr>() ||
NewD->hasAttr<CUDADeviceAttr>() || NewD->hasAttr<CUDAGlobalAttr>())
diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp
index 6ae825b9709..d12dda66dbf 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -3275,6 +3275,14 @@ ASTReader::ReadASTBlock(ModuleFile &F, unsigned ClientLoadCapabilities) {
UnusedLocalTypedefNameCandidates.push_back(
getGlobalDeclID(F, Record[I]));
break;
+
+ case CUDA_PRAGMA_FORCE_HOST_DEVICE_DEPTH:
+ if (Record.size() != 1) {
+ Error("invalid cuda pragma options record");
+ return Failure;
+ }
+ ForceCUDAHostDeviceDepth = Record[0];
+ break;
}
}
}
@@ -7128,6 +7136,7 @@ void ASTReader::UpdateSema() {
PragmaMSPointersToMembersState,
PointersToMembersPragmaLocation);
}
+ SemaObj->ForceCUDAHostDeviceDepth = ForceCUDAHostDeviceDepth;
}
IdentifierInfo *ASTReader::get(StringRef Name) {
diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
index cfe04433cfd..da513a7076e 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -1069,6 +1069,7 @@ void ASTWriter::WriteBlockInfoBlock() {
RECORD(POINTERS_TO_MEMBERS_PRAGMA_OPTIONS);
RECORD(UNUSED_LOCAL_TYPEDEF_NAME_CANDIDATES);
RECORD(DELETE_EXPRS_TO_ANALYZE);
+ RECORD(CUDA_PRAGMA_FORCE_HOST_DEVICE_DEPTH);
// SourceManager Block.
BLOCK(SOURCE_MANAGER_BLOCK);
@@ -3942,6 +3943,13 @@ void ASTWriter::WriteOpenCLExtensions(Sema &SemaRef) {
Stream.EmitRecord(OPENCL_EXTENSIONS, Record);
}
+void ASTWriter::WriteCUDAPragmas(Sema &SemaRef) {
+ if (SemaRef.ForceCUDAHostDeviceDepth > 0) {
+ RecordData::value_type Record[] = {SemaRef.ForceCUDAHostDeviceDepth};
+ Stream.EmitRecord(CUDA_PRAGMA_FORCE_HOST_DEVICE_DEPTH, Record);
+ }
+}
+
void ASTWriter::WriteObjCCategories() {
SmallVector<ObjCCategoriesInfo, 2> CategoriesMap;
RecordData Categories;
@@ -4619,6 +4627,7 @@ uint64_t ASTWriter::WriteASTCore(Sema &SemaRef, StringRef isysroot,
WriteIdentifierTable(PP, SemaRef.IdResolver, isModule);
WriteFPPragmaOptions(SemaRef.getFPOptions());
WriteOpenCLExtensions(SemaRef);
+ WriteCUDAPragmas(SemaRef);
WritePragmaDiagnosticMappings(Context.getDiagnostics(), isModule);
// If we're emitting a module, write out the submodule information.
OpenPOWER on IntegriCloud