[CUDA] Allow "extern __shared__ Foo foo[]" within anon. namespaces.
Summary:
Previously this triggered a -Wundefined-internal warning. But it's not
an undefined variable -- any variable of this form is a pointer to the
base of GPU core's shared memory.
Reviewers: tra
Subscribers: sanjoy, rsmith
Differential Revision: https://reviews.llvm.org/D46782
llvm-svn: 332621
diff --git a/clang/lib/AST/Decl.cpp b/clang/lib/AST/Decl.cpp
index 895d50d..08718e0 100644
--- a/clang/lib/AST/Decl.cpp
+++ b/clang/lib/AST/Decl.cpp
@@ -2432,6 +2432,23 @@
getASTContext().setTemplateOrSpecializationInfo(this, Template);
}
+bool VarDecl::isKnownToBeDefined() const {
+ const auto &LangOpts = getASTContext().getLangOpts();
+ // In CUDA mode without relocatable device code, variables of form 'extern
+ // __shared__ Foo foo[]' are pointers to the base of the GPU core's shared
+ // memory pool. These are never undefined variables, even if they appear
+ // inside of an anon namespace or static function.
+ //
+ // With CUDA relocatable device code enabled, these variables don't get
+ // special handling; they're treated like regular extern variables.
+ if (LangOpts.CUDA && !LangOpts.CUDARelocatableDeviceCode &&
+ hasExternalStorage() && hasAttr<CUDASharedAttr>() &&
+ isa<IncompleteArrayType>(getType()))
+ return true;
+
+ return hasDefinition();
+}
+
MemberSpecializationInfo *VarDecl::getMemberSpecializationInfo() const {
if (isStaticDataMember())
// FIXME: Remove ?