jlebar created this revision.
jlebar added a reviewer: rsmith.
jlebar added subscribers: cfe-commits, jhen, tra.
These cause us to consider all functions in-between to be __host__
__device__.
You can nest these pragmas; you just can't have more 'end's than
'begin's.
https://reviews.llvm.org/D24975
Files:
clang/include/clang/Basic/DiagnosticParseKinds.td
clang/include/clang/Parse/Parser.h
clang/include/clang/Sema/Sema.h
clang/lib/Parse/ParsePragma.cpp
clang/lib/Sema/SemaCUDA.cpp
clang/test/Parser/cuda-force-host-device.cu
Index: clang/test/Parser/cuda-force-host-device.cu
===================================================================
--- /dev/null
+++ clang/test/Parser/cuda-force-host-device.cu
@@ -0,0 +1,27 @@
+// RUN: %clang_cc1 -fsyntax-only -verify %s
+
+// Check the force_cuda_host_device_{begin,end} pragmas.
+
+#pragma clang force_cuda_host_device_begin
+void f();
+#pragma clang force_cuda_host_device_begin
+void g();
+#pragma clang force_cuda_host_device_end
+void h();
+#pragma clang force_cuda_host_device_end
+
+void i(); // expected-note {{not viable}}
+
+void host() {
+ f();
+ g();
+ h();
+ i();
+}
+
+__attribute__((device)) void device() {
+ f();
+ g();
+ h();
+ i(); // expected-error {{no matching function}}
+}
Index: clang/lib/Sema/SemaCUDA.cpp
===================================================================
--- clang/lib/Sema/SemaCUDA.cpp
+++ clang/lib/Sema/SemaCUDA.cpp
@@ -23,6 +23,19 @@
#include "llvm/ADT/SmallVector.h"
using namespace clang;
+void Sema::PushForceCUDAHostDevice() {
+ assert(getLangOpts().CUDA && "May be called only for CUDA compilations.");
+ ForceCUDAHostDeviceDepth++;
+}
+
+bool Sema::PopForceCUDAHostDevice() {
+ assert(getLangOpts().CUDA && "May be called only for CUDA compilations.");
+ if (ForceCUDAHostDeviceDepth == 0)
+ return false;
+ ForceCUDAHostDeviceDepth--;
+ return true;
+}
+
ExprResult Sema::ActOnCUDAExecConfigExpr(Scope *S, SourceLocation LLLLoc,
MultiExprArg ExecConfig,
SourceLocation GGGLoc) {
@@ -441,9 +454,23 @@
// * 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 && "May be called only for CUDA compilations.");
+
+ 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>())
Index: clang/lib/Parse/ParsePragma.cpp
===================================================================
--- clang/lib/Parse/ParsePragma.cpp
+++ clang/lib/Parse/ParsePragma.cpp
@@ -167,6 +167,26 @@
Token &FirstToken) override;
};
+struct PragmaForceCUDAHostDeviceStartHandler : public PragmaHandler {
+ PragmaForceCUDAHostDeviceStartHandler(Sema &Actions)
+ : PragmaHandler("force_cuda_host_device_begin"), Actions(Actions) {}
+ void HandlePragma(Preprocessor &PP, PragmaIntroducerKind Introducer,
+ Token &NameTok) override;
+
+private:
+ Sema &Actions;
+};
+
+struct PragmaForceCUDAHostDeviceEndHandler : public PragmaHandler {
+ PragmaForceCUDAHostDeviceEndHandler(Sema &Actions)
+ : PragmaHandler("force_cuda_host_device_end"), Actions(Actions) {}
+ void HandlePragma(Preprocessor &PP, PragmaIntroducerKind Introducer,
+ Token &NameTok) override;
+
+private:
+ Sema &Actions;
+};
+
} // end namespace
void Parser::initializePragmaHandlers() {
@@ -239,6 +259,15 @@
PP.AddPragmaHandler(MSIntrinsic.get());
}
+ if (getLangOpts().CUDA) {
+ CUDAForceHostDeviceStartHandler.reset(
+ new PragmaForceCUDAHostDeviceStartHandler(Actions));
+ PP.AddPragmaHandler("clang", CUDAForceHostDeviceStartHandler.get());
+ CUDAForceHostDeviceEndHandler.reset(
+ new PragmaForceCUDAHostDeviceEndHandler(Actions));
+ PP.AddPragmaHandler("clang", CUDAForceHostDeviceEndHandler.get());
+ }
+
OptimizeHandler.reset(new PragmaOptimizeHandler(Actions));
PP.AddPragmaHandler("clang", OptimizeHandler.get());
@@ -309,6 +338,13 @@
MSIntrinsic.reset();
}
+ if (getLangOpts().CUDA) {
+ PP.RemovePragmaHandler("clang", CUDAForceHostDeviceStartHandler.get());
+ CUDAForceHostDeviceStartHandler.reset();
+ PP.RemovePragmaHandler("clang", CUDAForceHostDeviceEndHandler.get());
+ CUDAForceHostDeviceEndHandler.reset();
+ }
+
PP.RemovePragmaHandler("STDC", FPContractHandler.get());
FPContractHandler.reset();
@@ -2187,3 +2223,13 @@
PP.Diag(Tok.getLocation(), diag::warn_pragma_extra_tokens_at_eol)
<< "intrinsic";
}
+void PragmaForceCUDAHostDeviceStartHandler::HandlePragma(
+ Preprocessor &PP, PragmaIntroducerKind Introducer, Token &NameTok) {
+ Actions.PushForceCUDAHostDevice();
+}
+void PragmaForceCUDAHostDeviceEndHandler::HandlePragma(
+ Preprocessor &PP, PragmaIntroducerKind Introducer, Token &NameTok) {
+ if (!Actions.PopForceCUDAHostDevice())
+ PP.Diag(NameTok.getLocation(),
+ diag::err_pragma_cannot_end_force_cuda_host_device);
+}
Index: clang/include/clang/Sema/Sema.h
===================================================================
--- clang/include/clang/Sema/Sema.h
+++ clang/include/clang/Sema/Sema.h
@@ -9185,6 +9185,20 @@
QualType FieldTy, bool IsMsStruct,
Expr *BitWidth, bool *ZeroWidth = nullptr);
+private:
+ unsigned ForceCUDAHostDeviceDepth = 0;
+
+public:
+ /// Increments our count of the number of times we've seen a pragma forcing
+ /// functions to be __host__ __device__. So long as this count is greater
+ /// than zero, all functions encountered will be __host__ __device__.
+ void PushForceCUDAHostDevice();
+
+ /// Decrements our count of the number of times we've seen a pragma forcing
+ /// functions to be __host__ __device__. Returns false if the count is 0
+ /// before incrementing, so you can emit an error.
+ bool PopForceCUDAHostDevice();
+
enum CUDAFunctionTarget {
CFT_Device,
CFT_Global,
Index: clang/include/clang/Parse/Parser.h
===================================================================
--- clang/include/clang/Parse/Parser.h
+++ clang/include/clang/Parse/Parser.h
@@ -173,6 +173,8 @@
std::unique_ptr<PragmaHandler> MSSection;
std::unique_ptr<PragmaHandler> MSRuntimeChecks;
std::unique_ptr<PragmaHandler> MSIntrinsic;
+ std::unique_ptr<PragmaHandler> CUDAForceHostDeviceStartHandler;
+ std::unique_ptr<PragmaHandler> CUDAForceHostDeviceEndHandler;
std::unique_ptr<PragmaHandler> OptimizeHandler;
std::unique_ptr<PragmaHandler> LoopHintHandler;
std::unique_ptr<PragmaHandler> UnrollHintHandler;
Index: clang/include/clang/Basic/DiagnosticParseKinds.td
===================================================================
--- clang/include/clang/Basic/DiagnosticParseKinds.td
+++ clang/include/clang/Basic/DiagnosticParseKinds.td
@@ -1022,6 +1022,10 @@
def warn_pragma_unroll_cuda_value_in_parens : Warning<
"argument to '#pragma unroll' should not be in parentheses in CUDA C/C++">,
InGroup<CudaCompat>;
+
+def err_pragma_cannot_end_force_cuda_host_device : Error<
+ "force_cuda_host_device_end pragma without matching "
+ "force_cuda_host_device_begin.">;
} // end of Parse Issue category.
let CategoryName = "Modules Issue" in {
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
