replica
/
hanchenye-llvm-project
2
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

[CUDA] Check initializers of instantiated template variables. 

We were already performing checks on non-template variables,
but the checks on templated ones were missing.

Differential Revision: https://reviews.llvm.org/D45231

llvm-svn: 334143
This commit is contained in:
Artem Belevich 2018-06-06 22:37:25 +00:00
parent df61be70b2
commit e9fa53a09b
5 changed files with 85 additions and 52 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

10
clang/include/clang/Sema/Sema.h
View File

@ -10166,6 +10166,16 @@ public:
bool isEmptyCudaConstructor(SourceLocation Loc, CXXConstructorDecl *CD);
bool isEmptyCudaDestructor(SourceLocation Loc, CXXDestructorDecl *CD);
// \brief Checks that initializers of \p Var satisfy CUDA restrictions. In
// case of error emits appropriate diagnostic and invalidates \p Var.
//
// \details CUDA allows only empty constructors as initializers for global
// variables (see E.2.3.1, CUDA 7.5). The same restriction also applies to all
// __shared__ variables whether they are local or not (they all are implicitly
// static in CUDA). One exception is that CUDA allows constant initializers
// for __constant__ and __device__ variables.
void checkAllowedCUDAInitializer(VarDecl *Var);
/// Check whether NewFD is a valid overload for CUDA. Emits
/// diagnostics and invalidates NewFD if not.
void checkCUDATargetOverload(FunctionDecl *NewFD,

53
clang/lib/Sema/SemaCUDA.cpp
View File

@ -472,6 +472,59 @@ bool Sema::isEmptyCudaDestructor(SourceLocation Loc, CXXDestructorDecl *DD) {
return true;
}
void Sema::checkAllowedCUDAInitializer(VarDecl *VD) {
if (VD->isInvalidDecl() || !VD->hasInit() || !VD->hasGlobalStorage())
return;
const Expr *Init = VD->getInit();
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 =
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();
}
} 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();
}
}
}
}
// With -fcuda-host-device-constexpr, an unattributed constexpr function is
// treated as implicitly __host__ __device__, unless:
// * it is a variadic function (device-side variadic functions are not

54
clang/lib/Sema/SemaDecl.cpp
View File

@ -11675,58 +11675,8 @@ void 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) {
const Expr *Init = VD->getInit();
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 =
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();
}
} 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();
}
}
}
}
}
if (getLangOpts().CUDA)
checkAllowedCUDAInitializer(VD);
// Grab the dllimport or dllexport attribute off of the VarDecl.
const InheritableAttr *DLLAttr = getDLLAttr(VD);

3
clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
View File

@ -4224,6 +4224,9 @@ void Sema::InstantiateVariableInitializer(
ActOnUninitializedDecl(Var);
}
if (getLangOpts().CUDA)
checkAllowedCUDAInitializer(Var);
}
/// Instantiate the definition of the given variable from its

17
clang/test/SemaCUDA/device-var-init.cu
View File

@ -225,3 +225,20 @@ inline __host__ __device__ void hd_emitted_host_only() {
static int x = 42; // no error on device because this is never codegen'ed there.
}
void call_hd_emitted_host_only() { hd_emitted_host_only(); }
// Verify that we also check field initializers in instantiated structs.
struct NontrivialInitializer {
__host__ __device__ NontrivialInitializer() : x(43) {}
int x;
};
template <typename T>
__global__ void bar() {
__shared__ T bad;
// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
}
void instantiate() {
bar<NontrivialInitializer><<<1, 1>>>();
// expected-note@-1 {{in instantiation of function template specialization 'bar<NontrivialInitializer>' requested here}}
}