Don't eagerly emit a global static merged with a local extern.

When we are visiting the extern declaration of 'i' in

static int i = 99;
int foo() {
  extern int i;
  return i;
}

We should not try to handle it as if it was an function static. That is, we
must consider the written storage class.

Fixing this then exposes that the assert in EmitGlobalVarDeclLValue and the
if leading to its call are not completely accurate. They were passing before
because the second decl was marked as having external storage. I changed them
to check the linkage, which I find easier to understand.

Last but not least, there is something strange going on with cuda and opencl.
My guess is that the linkage computation for these languages needs to be
audited, but I didn't want to change that in this patch so I just updated
the storage classes to keep the current behavior.

Thanks to Reed Kotler for reporting this.

llvm-svn: 170827

1 files changed, 7 insertions, 2 deletions

diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 96e911d17bc..ebfb4b1e78c 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -4284,8 +4284,10 @@ Sema::ActOnVariableDeclarator(Scope *S, Declarator &D, DeclContext *DC,
   if (getLangOpts().OpenCL) {
     // Set up the special work-group-local storage class for variables in the
     // OpenCL __local address space.
-    if (R.getAddressSpace() == LangAS::opencl_local)
+    if (R.getAddressSpace() == LangAS::opencl_local) {
       SC = SC_OpenCLWorkGroupLocal;
+      SCAsWritten = SC_OpenCLWorkGroupLocal;
+    }
   }
 
   bool isExplicitSpecialization = false;
@@ -4420,8 +4422,11 @@ Sema::ActOnVariableDeclarator(Scope *S, Declarator &D, DeclContext *DC,
     // CUDA B.2.5: "__shared__ and __constant__ variables have implied static
     // storage [duration]."
     if (SC == SC_None && S->getFnParent() != 0 &&
-       (NewVD->hasAttr<CUDASharedAttr>() || NewVD->hasAttr<CUDAConstantAttr>()))
+        (NewVD->hasAttr<CUDASharedAttr>() ||
+         NewVD->hasAttr<CUDAConstantAttr>())) {
       NewVD->setStorageClass(SC_Static);
+      NewVD->setStorageClassAsWritten(SC_Static);
+    }
   }
 
   // In auto-retain/release, infer strong retension for variables of