Index: cfe/trunk/lib/CodeGen/CGDecl.cpp
===================================================================
--- cfe/trunk/lib/CodeGen/CGDecl.cpp
+++ cfe/trunk/lib/CodeGen/CGDecl.cpp
@@ -371,8 +371,15 @@
 
   llvm::GlobalVariable *var =
     cast<llvm::GlobalVariable>(addr->stripPointerCasts());
+
+  // CUDA's local and local static __shared__ variables should not
+  // have any non-empty initializers. This is ensured by Sema.
+  // Whatever initializer such variable may have when it gets here is
+  // a no-op and should not be emitted.
+  bool isCudaSharedVar = getLangOpts().CUDA && getLangOpts().CUDAIsDevice &&
+                         D.hasAttr<CUDASharedAttr>();
   // If this value has an initializer, emit it.
-  if (D.getInit())
+  if (D.getInit() && !isCudaSharedVar)
     var = AddInitializerToStaticVarDecl(D, var);
 
   var->setAlignment(alignment.getQuantity());
@@ -1874,4 +1881,3 @@
     return;
   getOpenMPRuntime().emitUserDefinedReduction(CGF, D);
 }
-
Index: cfe/trunk/lib/Sema/SemaDecl.cpp
===================================================================
--- cfe/trunk/lib/Sema/SemaDecl.cpp
+++ cfe/trunk/lib/Sema/SemaDecl.cpp
@@ -10414,14 +10414,15 @@
 
   // Perform check for initializers of device-side global variables.
   // CUDA allows empty constructors as initializers (see E.2.3.1, CUDA
-  // 7.5). CUDA also allows constant initializers for __constant__ and
-  // __device__ variables.
+  // 7.5). We must also apply the same checks to all __shared__
+  // variables whether they are local or not. CUDA also allows
+  // constant initializers for __constant__ and __device__ variables.
   if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice) {
     const Expr *Init = VD->getInit();
-    const bool IsGlobal = VD->hasGlobalStorage() && !VD->isStaticLocal();
-    if (Init && IsGlobal &&
+    if (Init && VD->hasGlobalStorage() &&
         (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 =
Index: cfe/trunk/test/CodeGenCUDA/address-spaces.cu
===================================================================
--- cfe/trunk/test/CodeGenCUDA/address-spaces.cu
+++ cfe/trunk/test/CodeGenCUDA/address-spaces.cu
@@ -25,8 +25,6 @@
 // CHECK: @_ZZ5func3vE1a = internal addrspace(3) global float 0.000000e+00
 // CHECK: @_ZZ5func4vE1a = internal addrspace(3) global float 0.000000e+00
 // CHECK: @b = addrspace(3) global float undef
-// CHECK: @c = addrspace(3) global %struct.c undef
-// CHECK  @d = addrspace(3) global %struct.d undef
 
 __device__ void foo() {
   // CHECK: load i32, i32* addrspacecast (i32 addrspace(1)* @i to i32*)
@@ -94,32 +92,3 @@
 }
 // CHECK: define float* @_Z5func5v()
 // CHECK: ret float* addrspacecast (float addrspace(3)* @b to float*)
-
-struct StructWithCtor {
-  __device__ StructWithCtor(): data(1) {}
-  __device__ StructWithCtor(const StructWithCtor &second): data(second.data) {}
-  __device__ int getData() { return data; }
-  int data;
-};
-
-__device__ int construct_shared_struct() {
-// CHECK-LABEL: define i32 @_Z23construct_shared_structv()
-  __shared__ StructWithCtor s;
-// CHECK: call void @_ZN14StructWithCtorC1Ev(%struct.StructWithCtor* addrspacecast (%struct.StructWithCtor addrspace(3)* @_ZZ23construct_shared_structvE1s to %struct.StructWithCtor*))
-  __shared__ StructWithCtor t(s);
-// CHECK: call void @_ZN14StructWithCtorC1ERKS_(%struct.StructWithCtor* addrspacecast (%struct.StructWithCtor addrspace(3)* @_ZZ23construct_shared_structvE1t to %struct.StructWithCtor*), %struct.StructWithCtor* dereferenceable(4) addrspacecast (%struct.StructWithCtor addrspace(3)* @_ZZ23construct_shared_structvE1s to %struct.StructWithCtor*))
-  return t.getData();
-// CHECK: call i32 @_ZN14StructWithCtor7getDataEv(%struct.StructWithCtor* addrspacecast (%struct.StructWithCtor addrspace(3)* @_ZZ23construct_shared_structvE1t to %struct.StructWithCtor*))
-}
-
-// Make sure we allow __shared__ structures with default or empty constructors.
-struct c {
-  int i;
-};
-__shared__ struct c c;
-
-struct d {
-  int i;
-  d() {}
-};
-__shared__ struct d d;
Index: cfe/trunk/test/CodeGenCUDA/device-var-init.cu
===================================================================
--- cfe/trunk/test/CodeGenCUDA/device-var-init.cu
+++ cfe/trunk/test/CodeGenCUDA/device-var-init.cu
@@ -63,6 +63,8 @@
 
 // static in-class field initializer.  NVCC does not allow it, but
 // clang generates static initializer for this, so we'll accept it.
+// We still can't use it on __shared__ vars as they don't allow *any*
+// initializers.
 struct NCFS {
   int ncfs = 3;
 };
@@ -367,8 +369,13 @@
   T_B_NEC t_b_nec;
   T_F_NEC t_f_nec;
   T_FA_NEC t_fa_nec;
-  static __shared__ UC s_uc;
+  static __shared__ EC s_ec;
+  static __shared__ ETC s_etc;
 #if ERROR_CASE
+  static __shared__ NCFS s_ncfs;
+  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+  static __shared__ UC s_uc;
+  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
   static __device__ int ds;
   // expected-error@-1 {{Within a __device__/__global__ function, only __shared__ variables may be marked "static"}}
   static __constant__ int dc;
@@ -394,7 +401,8 @@
 // CHECK:   call void @_ZN7T_B_NECC1Ev(%struct.T_B_NEC* %t_b_nec)
 // CHECK:   call void @_ZN7T_F_NECC1Ev(%struct.T_F_NEC* %t_f_nec)
 // CHECK:   call void @_ZN8T_FA_NECC1Ev(%struct.T_FA_NEC* %t_fa_nec)
-// CHECK:   call void @_ZN2UCC1Ev(%struct.UC* addrspacecast (%struct.UC addrspace(3)* @_ZZ2dfvE4s_uc to %struct.UC*))
+// CHECK-NOT: call void @_ZN2ECC1Ev(%struct.EC* addrspacecast (%struct.EC addrspace(3)* @_ZZ2dfvE4s_ec to %struct.EC*))
+// CHECK-NOT: call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* addrspacecast (%struct.ETC addrspace(3)* @_ZZ2dfvE5s_etc to %struct.ETC*))
 // CHECK: ret void
 
 // We should not emit global init function.