Author: tra Date: Wed Jun 6 15:37:25 2018 New Revision: 334143 URL: http://llvm.org/viewvc/llvm-project?rev=334143&view=rev Log: [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 Modified: cfe/trunk/include/clang/Sema/Sema.h cfe/trunk/lib/Sema/SemaCUDA.cpp cfe/trunk/lib/Sema/SemaDecl.cpp cfe/trunk/lib/Sema/SemaTemplateInstantiateDecl.cpp cfe/trunk/test/SemaCUDA/device-var-init.cu Modified: cfe/trunk/include/clang/Sema/Sema.h URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Sema/Sema.h?rev=334143&r1=334142&r2=334143&view=diff ============================================================================== --- cfe/trunk/include/clang/Sema/Sema.h (original) +++ cfe/trunk/include/clang/Sema/Sema.h Wed Jun 6 15:37:25 2018 @@ -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, Modified: cfe/trunk/lib/Sema/SemaCUDA.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaCUDA.cpp?rev=334143&r1=334142&r2=334143&view=diff ============================================================================== --- cfe/trunk/lib/Sema/SemaCUDA.cpp (original) +++ cfe/trunk/lib/Sema/SemaCUDA.cpp Wed Jun 6 15:37:25 2018 @@ -472,6 +472,59 @@ bool Sema::isEmptyCudaDestructor(SourceL 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 Modified: cfe/trunk/lib/Sema/SemaDecl.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDecl.cpp?rev=334143&r1=334142&r2=334143&view=diff ============================================================================== --- cfe/trunk/lib/Sema/SemaDecl.cpp (original) +++ cfe/trunk/lib/Sema/SemaDecl.cpp Wed Jun 6 15:37:25 2018 @@ -11675,58 +11675,8 @@ void Sema::FinalizeDeclaration(Decl *Thi // 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); Modified: cfe/trunk/lib/Sema/SemaTemplateInstantiateDecl.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaTemplateInstantiateDecl.cpp?rev=334143&r1=334142&r2=334143&view=diff ============================================================================== --- cfe/trunk/lib/Sema/SemaTemplateInstantiateDecl.cpp (original) +++ cfe/trunk/lib/Sema/SemaTemplateInstantiateDecl.cpp Wed Jun 6 15:37:25 2018 @@ -4224,6 +4224,9 @@ void Sema::InstantiateVariableInitialize ActOnUninitializedDecl(Var); } + + if (getLangOpts().CUDA) + checkAllowedCUDAInitializer(Var); } /// Instantiate the definition of the given variable from its Modified: cfe/trunk/test/SemaCUDA/device-var-init.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/device-var-init.cu?rev=334143&r1=334142&r2=334143&view=diff ============================================================================== --- cfe/trunk/test/SemaCUDA/device-var-init.cu (original) +++ cfe/trunk/test/SemaCUDA/device-var-init.cu Wed Jun 6 15:37:25 2018 @@ -225,3 +225,20 @@ inline __host__ __device__ void hd_emitt 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}} +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
