[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

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@332621 91177308-0d34-0410-b5e6-96231b3b80d8

diff --git a/include/clang/AST/Decl.h b/include/clang/AST/Decl.h
index d16be216807026cc7d50887f950ffbec9bd74670..c10ca9e9670b7d327137c8d2d8416df9c3ac3659 100644 (file)
--- a/include/clang/AST/Decl.h
+++ b/include/clang/AST/Decl.h
@@ -1456,6 +1456,11 @@ public:
 
   void setDescribedVarTemplate(VarTemplateDecl *Template);
 
+  // Is this variable known to have a definition somewhere in the complete
+  // program? This may be true even if the declaration has internal linkage and
+  // has no definition within this source file.
+  bool isKnownToBeDefined() const;
+
   // Implement isa/cast/dyncast/etc.
   static bool classof(const Decl *D) { return classofKind(D->getKind()); }
   static bool classofKind(Kind K) { return K >= firstVar && K <= lastVar; }

diff --git a/lib/AST/Decl.cpp b/lib/AST/Decl.cpp
index 895d50d1013487add9032ea98cc165551d5f6ea5..08718e0a9e8f70260101081e63dae835b29a5ed0 100644 (file)
--- a/lib/AST/Decl.cpp
+++ b/lib/AST/Decl.cpp
@@ -2432,6 +2432,23 @@ void VarDecl::setDescribedVarTemplate(VarTemplateDecl *Template) {
   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 ?

diff --git a/lib/Sema/Sema.cpp b/lib/Sema/Sema.cpp
index c0997688b2d022ea635088a96c796ea3e0a23fd6..aa4f6461b65e60e1bb0244cfd6ce61b938f6530d 100644 (file)
--- a/lib/Sema/Sema.cpp
+++ b/lib/Sema/Sema.cpp
@@ -653,6 +653,11 @@ void Sema::getUndefinedButUsed(
           !isExternalWithNoLinkageType(VD) &&
           !VD->getMostRecentDecl()->isInline())
         continue;
+
+      // Skip VarDecls that lack formal definitions but which we know are in
+      // fact defined somewhere.
+      if (VD->isKnownToBeDefined())
+        continue;
     }
 
     Undefined.push_back(std::make_pair(ND, UndefinedUse.second));

diff --git a/test/SemaCUDA/extern-shared.cu b/test/SemaCUDA/extern-shared.cu
index 2c0c6e0e3df7873f7c5b9c40f875714cfd7748c8..ec6f31695052a233073c34a44aa96c1bcd902160 100644 (file)
--- a/test/SemaCUDA/extern-shared.cu
+++ b/test/SemaCUDA/extern-shared.cu
@@ -1,10 +1,10 @@
-// RUN: %clang_cc1 -fsyntax-only -verify %s
-// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s
+// RUN: %clang_cc1 -fsyntax-only -Wundefined-internal -verify %s
+// RUN: %clang_cc1 -fsyntax-only -Wundefined-internal -fcuda-is-device -verify %s
 
-// RUN: %clang_cc1 -fsyntax-only -fcuda-rdc -verify=rdc %s
-// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -fcuda-rdc -verify=rdc %s
-// These declarations are fine in separate compilation mode:
-// rdc-no-diagnostics
+// RUN: %clang_cc1 -fsyntax-only -Wundefined-internal -fcuda-rdc -verify=rdc %s
+// RUN: %clang_cc1 -fsyntax-only -Wundefined-internal -fcuda-is-device -fcuda-rdc -verify=rdc %s
+
+// Most of these declarations are fine in separate compilation mode.
 
 #include "Inputs/cuda.h"
 
@@ -26,3 +26,18 @@ __host__ __device__ void bar() {
 extern __shared__ int global; // expected-error {{__shared__ variable 'global' cannot be 'extern'}}
 extern __shared__ int global_arr[]; // ok
 extern __shared__ int global_arr1[1]; // expected-error {{__shared__ variable 'global_arr1' cannot be 'extern'}}
+
+// Check that, iff we're not in rdc mode, extern __shared__ can appear in an
+// anonymous namespace / in a static function without generating a warning
+// about a variable with internal linkage but no definition
+// (-Wundefined-internal).
+namespace {
+extern __shared__ int global_arr[]; // rdc-warning {{has internal linkage but is not defined}}
+__global__ void in_anon_ns() {
+  extern __shared__ int local_arr[]; // rdc-warning {{has internal linkage but is not defined}}
+
+  // Touch arrays to generate the warning.
+  local_arr[0] = 0;  // rdc-note {{used here}}
+  global_arr[0] = 0; // rdc-note {{used here}}
+}
+} // namespace