r278196 - [CUDA] Reject calls to __device__ functions from host variable global initializers.
Justin Lebar via cfe-commits
cfe-commits at lists.llvm.org
Tue Aug 9 18:09:22 PDT 2016
Author: jlebar
Date: Tue Aug 9 20:09:21 2016
New Revision: 278196
URL: http://llvm.org/viewvc/llvm-project?rev=278196&view=rev
Log:
[CUDA] Reject calls to __device__ functions from host variable global initializers.
Reviewers: tra
Subscribers: cfe-commits
Differential Revision: https://reviews.llvm.org/D23335
Added:
cfe/trunk/test/SemaCUDA/global-initializers-host.cu
Modified:
cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
cfe/trunk/lib/Sema/SemaDecl.cpp
Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td?rev=278196&r1=278195&r2=278196&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td (original)
+++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td Tue Aug 9 20:09:21 2016
@@ -6640,6 +6640,9 @@ def err_global_call_not_config : Error<
def err_ref_bad_target : Error<
"reference to %select{__device__|__global__|__host__|__host__ __device__}0 "
"function %1 in %select{__device__|__global__|__host__|__host__ __device__}2 function">;
+def err_ref_bad_target_global_initializer : Error<
+ "reference to %select{__device__|__global__|__host__|__host__ __device__}0 "
+ "function %1 in global initializer">;
def warn_kern_is_method : Extension<
"kernel function %0 is a member function; this may not be accepted by nvcc">,
InGroup<CudaCompat>;
Modified: cfe/trunk/lib/Sema/SemaDecl.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDecl.cpp?rev=278196&r1=278195&r2=278196&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaDecl.cpp (original)
+++ cfe/trunk/lib/Sema/SemaDecl.cpp Tue Aug 9 20:09:21 2016
@@ -10728,36 +10728,55 @@ Sema::FinalizeDeclaration(Decl *ThisDecl
// 7.5). We must also apply the same checks to all __shared__
// variables whether they are local or not. CUDA also allows
// constant initializers for __constant__ and __device__ variables.
- if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice) {
+ if (getLangOpts().CUDA) {
const Expr *Init = VD->getInit();
- if (Init && VD->hasGlobalStorage() &&
- (VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>() ||
- VD->hasAttr<CUDASharedAttr>())) {
- assert((!VD->isStaticLocal() || VD->hasAttr<CUDASharedAttr>()));
- bool AllowedInit = false;
- if (const CXXConstructExpr *CE = dyn_cast<CXXConstructExpr>(Init))
- AllowedInit =
- isEmptyCudaConstructor(VD->getLocation(), CE->getConstructor());
- // We'll allow constant initializers even if it's a non-empty
- // constructor according to CUDA rules. This deviates from NVCC,
- // but allows us to handle things like constexpr constructors.
- if (!AllowedInit &&
- (VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>()))
- AllowedInit = VD->getInit()->isConstantInitializer(
- Context, VD->getType()->isReferenceType());
-
- // Also make sure that destructor, if there is one, is empty.
- if (AllowedInit)
- if (CXXRecordDecl *RD = VD->getType()->getAsCXXRecordDecl())
+ if (Init && VD->hasGlobalStorage()) {
+ if (VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>() ||
+ VD->hasAttr<CUDASharedAttr>()) {
+ assert((!VD->isStaticLocal() || VD->hasAttr<CUDASharedAttr>()));
+ bool AllowedInit = false;
+ if (const CXXConstructExpr *CE = dyn_cast<CXXConstructExpr>(Init))
AllowedInit =
- isEmptyCudaDestructor(VD->getLocation(), RD->getDestructor());
+ isEmptyCudaConstructor(VD->getLocation(), CE->getConstructor());
+ // We'll allow constant initializers even if it's a non-empty
+ // constructor according to CUDA rules. This deviates from NVCC,
+ // but allows us to handle things like constexpr constructors.
+ if (!AllowedInit &&
+ (VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>()))
+ AllowedInit = VD->getInit()->isConstantInitializer(
+ Context, VD->getType()->isReferenceType());
+
+ // Also make sure that destructor, if there is one, is empty.
+ if (AllowedInit)
+ if (CXXRecordDecl *RD = VD->getType()->getAsCXXRecordDecl())
+ AllowedInit =
+ isEmptyCudaDestructor(VD->getLocation(), RD->getDestructor());
- if (!AllowedInit) {
- Diag(VD->getLocation(), VD->hasAttr<CUDASharedAttr>()
- ? diag::err_shared_var_init
- : diag::err_dynamic_var_init)
- << Init->getSourceRange();
- VD->setInvalidDecl();
+ if (!AllowedInit) {
+ Diag(VD->getLocation(), VD->hasAttr<CUDASharedAttr>()
+ ? diag::err_shared_var_init
+ : diag::err_dynamic_var_init)
+ << Init->getSourceRange();
+ VD->setInvalidDecl();
+ }
+ } else {
+ // This is a host-side global variable. Check that the initializer is
+ // callable from the host side.
+ const FunctionDecl *InitFn = nullptr;
+ if (const CXXConstructExpr *CE = dyn_cast<CXXConstructExpr>(Init)) {
+ InitFn = CE->getConstructor();
+ } else if (const CallExpr *CE = dyn_cast<CallExpr>(Init)) {
+ InitFn = CE->getDirectCallee();
+ }
+ if (InitFn) {
+ CUDAFunctionTarget InitFnTarget = IdentifyCUDATarget(InitFn);
+ if (InitFnTarget != CFT_Host && InitFnTarget != CFT_HostDevice) {
+ Diag(VD->getLocation(), diag::err_ref_bad_target_global_initializer)
+ << InitFnTarget << InitFn;
+ Diag(InitFn->getLocation(), diag::note_previous_decl) << InitFn;
+ VD->setInvalidDecl();
+ }
+ }
}
}
}
Added: cfe/trunk/test/SemaCUDA/global-initializers-host.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/global-initializers-host.cu?rev=278196&view=auto
==============================================================================
--- cfe/trunk/test/SemaCUDA/global-initializers-host.cu (added)
+++ cfe/trunk/test/SemaCUDA/global-initializers-host.cu Tue Aug 9 20:09:21 2016
@@ -0,0 +1,32 @@
+// RUN: %clang_cc1 %s --std=c++11 -triple x86_64-linux-unknown -fsyntax-only -o - -verify
+
+#include "Inputs/cuda.h"
+
+// Check that we get an error if we try to call a __device__ function from a
+// module initializer.
+
+struct S {
+ __device__ S() {}
+ // expected-note at -1 {{'S' declared here}}
+};
+
+S s;
+// expected-error at -1 {{reference to __device__ function 'S' in global initializer}}
+
+struct T {
+ __host__ __device__ T() {}
+};
+T t; // No error, this is OK.
+
+struct U {
+ __host__ U() {}
+ __device__ U(int) {}
+ // expected-note at -1 {{'U' declared here}}
+};
+U u(42);
+// expected-error at -1 {{reference to __device__ function 'U' in global initializer}}
+
+__device__ int device_fn() { return 42; }
+// expected-note at -1 {{'device_fn' declared here}}
+int n = device_fn();
+// expected-error at -1 {{reference to __device__ function 'device_fn' in global initializer}}
More information about the cfe-commits
mailing list