jlebar updated this revision to Diff 72717.
jlebar marked 2 inline comments as done.
jlebar added a comment.
Address Richard Smith's review comments:
- Change macro format.
- Add tests (these Just Worked).
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-templates.cu
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,33 @@
+// RUN: %clang_cc1 -fsyntax-only -verify %s
+
+// Check the force_cuda_host_device pragma.
+
+#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}}
+}
+
+#pragma clang force_cuda_host_device foo
+// expected-warning@-1 {{Incorrect use of #pragma clang force_cuda_host_device begin|end}}
+
+#pragma clang force_cuda_host_device
+// expected-warning@-1 {{Incorrect use of #pragma clang force_cuda_host_device begin|end}}
Index: clang/test/Parser/cuda-force-host-device-templates.cu
===================================================================
--- /dev/null
+++ clang/test/Parser/cuda-force-host-device-templates.cu
@@ -0,0 +1,41 @@
+// RUN: %clang_cc1 -S -verify -fcuda-is-device %s -o /dev/null
+
+// Check how the force_cuda_host_device pragma interacts with template
+// instantiations. The errors here are emitted at codegen, so we can't do
+// -fsyntax-only.
+
+template <typename T>
+T foo() { // expected-note {{declared here}}
+ return T();
+}
+
+template <typename T>
+struct X {
+ void foo(); // expected-note {{declared here}}
+};
+
+#pragma clang force_cuda_host_device begin
+__attribute__((host)) __attribute__((device)) void test() {
+ int n = foo<int>(); // expected-error {{reference to __host__ function 'foo<int>'}}
+ X<int>().foo(); // expected-error {{reference to __host__ function 'foo'}}
+}
+#pragma clang force_cuda_host_device end
+
+// Same thing as above, but within a force_cuda_host_device block without a
+// corresponding end.
+
+template <typename T>
+T bar() { // expected-note {{declared here}}
+ return T();
+}
+
+template <typename T>
+struct Y {
+ void bar(); // expected-note {{declared here}}
+};
+
+#pragma clang force_cuda_host_device begin
+__attribute__((host)) __attribute__((device)) void test2() {
+ int n = bar<int>(); // expected-error {{reference to __host__ function 'bar<int>'}}
+ Y<int>().bar(); // expected-error {{reference to __host__ function 'bar'}}
+}
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,16 @@
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 @@
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 @@
MSIntrinsic.reset();
}
+ if (getLangOpts().CUDA) {
+ PP.RemovePragmaHandler("clang", CUDAForceHostDeviceHandler.get());
+ CUDAForceHostDeviceHandler.reset();
+ }
+
PP.RemovePragmaHandler("STDC", FPContractHandler.get());
FPContractHandler.reset();
@@ -2187,3 +2208,19 @@
PP.Diag(Tok.getLocation(), diag::warn_pragma_extra_tokens_at_eol)
<< "intrinsic";
}
+void PragmaForceCUDAHostDeviceHandler::HandlePragma(
+ Preprocessor &PP, PragmaIntroducerKind Introducer, Token &Tok) {
+ auto FirstTok = Tok;
+
+ PP.Lex(Tok);
+ IdentifierInfo *Info = Tok.getIdentifierInfo();
+ if (Info && Info->isStr("begin")) {
+ Actions.PushForceCUDAHostDevice();
+ } else if (Info && Info->isStr("end")) {
+ if (!Actions.PopForceCUDAHostDevice())
+ PP.Diag(FirstTok.getLocation(),
+ diag::err_pragma_cannot_end_force_cuda_host_device);
+ } else
+ PP.Diag(FirstTok.getLocation(),
+ diag::warn_pragma_force_cuda_host_device_bad_arg);
+}
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,7 @@
std::unique_ptr<PragmaHandler> MSSection;
std::unique_ptr<PragmaHandler> MSRuntimeChecks;
std::unique_ptr<PragmaHandler> MSIntrinsic;
+ std::unique_ptr<PragmaHandler> CUDAForceHostDeviceHandler;
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,13 @@
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 warn_pragma_force_cuda_host_device_bad_arg : Warning<
+ "Incorrect use of #pragma clang force_cuda_host_device begin|end">;
+
+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
