diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -426,6 +426,10 @@
   if (CD->getParent()->isDynamicClass())
     return false;
 
+  // Union ctor does not call ctors of its data members.
+  if (CD->getParent()->isUnion())
+    return true;
+
   // The only form of initializer allowed is an empty constructor.
   // This will recursively check all base classes and member initializers
   if (!llvm::all_of(CD->inits(), [&](const CXXCtorInitializer *CI) {
@@ -465,6 +469,11 @@
   if (ClassDecl->isDynamicClass())
     return false;
 
+  // Union does not have base class and union dtor does not call dtors of its
+  // data members.
+  if (DD->getParent()->isUnion())
+    return true;
+
   // Only empty destructors are allowed. This will recursively check
   // destructors for all base classes...
   if (!llvm::all_of(ClassDecl->bases(), [&](const CXXBaseSpecifier &BS) {
diff --git a/clang/test/SemaCUDA/union-init.cu b/clang/test/SemaCUDA/union-init.cu
new file mode 100644
--- /dev/null
+++ b/clang/test/SemaCUDA/union-init.cu
@@ -0,0 +1,44 @@
+// RUN: %clang_cc1 %s --std=c++11 -triple x86_64-linux-unknown -fsyntax-only -o - -verify
+
+#include "Inputs/cuda.h"
+
+struct A {
+  int a;
+  __device__ A() { a = 1; }
+  __device__ ~A() { a = 2; }
+};
+
+// This can be a global var since ctor/dtors of data members are not called.
+union B {
+  A a;
+  __device__ B() {}
+  __device__ ~B() {}
+};
+
+// This cannot be a global var since it has a dynamic ctor.
+union C {
+  A a;
+  __device__ C() { a.a = 3; }
+  __device__ ~C() {}
+};
+
+// This cannot be a global var since it has a dynamic dtor.
+union D {
+  A a;
+  __device__ D() { }
+  __device__ ~D() { a.a = 4; }
+};
+
+__device__ B b;
+__device__ C c;
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__device__ D d;
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+__device__ void foo() {
+  __shared__ B b;
+  __shared__ C c;
+  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+  __shared__ D d;
+  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+}