Index: clang/lib/Sema/Sema.cpp
===================================================================
--- clang/lib/Sema/Sema.cpp
+++ clang/lib/Sema/Sema.cpp
@@ -653,6 +653,15 @@
           !isExternalWithNoLinkageType(VD) &&
           !VD->getMostRecentDecl()->isInline())
         continue;
+
+      // In CUDA mode, function-local 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.
+      if (getLangOpts().CUDA && VD->isLocalExternDecl() &&
+          VD->hasAttr<CUDASharedAttr>() && VD->getType()->isArrayType()) {
+        continue;
+      }
     }
 
     Undefined.push_back(std::make_pair(ND, UndefinedUse.second));
Index: clang/test/SemaCUDA/extern-shared.cu
===================================================================
--- clang/test/SemaCUDA/extern-shared.cu
+++ clang/test/SemaCUDA/extern-shared.cu
@@ -1,8 +1,8 @@
-// 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
+// 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
 // These declarations are fine in separate compilation mode:
 // rdc-no-diagnostics
 
@@ -26,3 +26,14 @@
 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 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 {
+__global__ void in_anon_ns() {
+  extern __shared__ int arr[]; // ok, no warning about having internal linkage but not being defined.
+  // Touch arr to generate the warning.
+  arr[0] = 0;
+}
+} // namespace