aboutsummaryrefslogtreecommitdiffstats
path: root/lib/Sema/SemaCUDA.cpp
diff options
context:
space:
mode:
authorDimitry Andric <dim@FreeBSD.org>2018-07-28 11:06:01 +0000
committerDimitry Andric <dim@FreeBSD.org>2018-07-28 11:06:01 +0000
commit486754660bb926339aefcf012a3f848592babb8b (patch)
treeecdbc446c9876f4f120f701c243373cd3cb43db3 /lib/Sema/SemaCUDA.cpp
parent55e6d896ad333f07bb3b1ba487df214fc268a4ab (diff)
downloadsrc-486754660bb926339aefcf012a3f848592babb8b.tar.gz
src-486754660bb926339aefcf012a3f848592babb8b.zip
Vendor import of clang trunk r338150:vendor/clang/clang-trunk-r338150
Notes
Notes: svn path=/vendor/clang/dist/; revision=336815 svn path=/vendor/clang/clang-trunk-r338150/; revision=336816; tag=vendor/clang/clang-trunk-r338150
Diffstat (limited to 'lib/Sema/SemaCUDA.cpp')
-rw-r--r--lib/Sema/SemaCUDA.cpp88
1 files changed, 73 insertions, 15 deletions
diff --git a/lib/Sema/SemaCUDA.cpp b/lib/Sema/SemaCUDA.cpp
index cac5f682275e..13dd8d936fd2 100644
--- a/lib/Sema/SemaCUDA.cpp
+++ b/lib/Sema/SemaCUDA.cpp
@@ -7,7 +7,7 @@
//
//===----------------------------------------------------------------------===//
/// \file
-/// \brief This file implements semantic analysis for CUDA constructs.
+/// This file implements semantic analysis for CUDA constructs.
///
//===----------------------------------------------------------------------===//
@@ -42,8 +42,9 @@ ExprResult Sema::ActOnCUDAExecConfigExpr(Scope *S, SourceLocation LLLLoc,
SourceLocation GGGLoc) {
FunctionDecl *ConfigDecl = Context.getcudaConfigureCallDecl();
if (!ConfigDecl)
- return ExprError(Diag(LLLLoc, diag::err_undeclared_var_use)
- << "cudaConfigureCall");
+ return ExprError(
+ Diag(LLLLoc, diag::err_undeclared_var_use)
+ << (getLangOpts().HIP ? "hipConfigureCall" : "cudaConfigureCall"));
QualType ConfigQTy = ConfigDecl->getType();
DeclRefExpr *ConfigDR = new (Context)
@@ -54,30 +55,31 @@ ExprResult Sema::ActOnCUDAExecConfigExpr(Scope *S, SourceLocation LLLLoc,
/*IsExecConfig=*/true);
}
-Sema::CUDAFunctionTarget Sema::IdentifyCUDATarget(const AttributeList *Attr) {
+Sema::CUDAFunctionTarget
+Sema::IdentifyCUDATarget(const ParsedAttributesView &Attrs) {
bool HasHostAttr = false;
bool HasDeviceAttr = false;
bool HasGlobalAttr = false;
bool HasInvalidTargetAttr = false;
- while (Attr) {
- switch(Attr->getKind()){
- case AttributeList::AT_CUDAGlobal:
+ for (const ParsedAttr &AL : Attrs) {
+ switch (AL.getKind()) {
+ case ParsedAttr::AT_CUDAGlobal:
HasGlobalAttr = true;
break;
- case AttributeList::AT_CUDAHost:
+ case ParsedAttr::AT_CUDAHost:
HasHostAttr = true;
break;
- case AttributeList::AT_CUDADevice:
+ case ParsedAttr::AT_CUDADevice:
HasDeviceAttr = true;
break;
- case AttributeList::AT_CUDAInvalidTarget:
+ case ParsedAttr::AT_CUDAInvalidTarget:
HasInvalidTargetAttr = true;
break;
default:
break;
}
- Attr = Attr->getNext();
}
+
if (HasInvalidTargetAttr)
return CFT_InvalidTarget;
@@ -471,6 +473,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
@@ -521,7 +576,7 @@ void Sema::maybeAddCUDAHostDeviceAttrs(FunctionDecl *NewD,
if (!getSourceManager().isInSystemHeader(Match->getLocation())) {
Diag(NewD->getLocation(),
diag::err_cuda_unattributed_constexpr_cannot_overload_device)
- << NewD->getName();
+ << NewD;
Diag(Match->getLocation(),
diag::note_cuda_conflicting_device_function_declared_here);
}
@@ -790,9 +845,12 @@ bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) {
// If the caller is known-emitted, mark the callee as known-emitted.
// Otherwise, mark the call in our call graph so we can traverse it later.
bool CallerKnownEmitted = IsKnownEmitted(*this, Caller);
- if (CallerKnownEmitted)
- MarkKnownEmitted(*this, Caller, Callee, Loc);
- else {
+ if (CallerKnownEmitted) {
+ // Host-side references to a __global__ function refer to the stub, so the
+ // function itself is never emitted and therefore should not be marked.
+ if (getLangOpts().CUDAIsDevice || IdentifyCUDATarget(Callee) != CFT_Global)
+ MarkKnownEmitted(*this, Caller, Callee, Loc);
+ } else {
// If we have
// host fn calls kernel fn calls host+device,
// the HD function does not get instantiated on the host. We model this by