diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp
index f441bf4eda1b..955b86a38b26 100644
--- a/clang/lib/CodeGen/CGCall.cpp
+++ b/clang/lib/CodeGen/CGCall.cpp
@@ -1814,6 +1814,9 @@ void CodeGenModule::ConstructAttributeList(
     // them).  LLVM will remove this attribute where it safely can.
     FuncAttrs.addAttribute(llvm::Attribute::Convergent);
 
+    // Exceptions aren't supported in CUDA device code.
+    FuncAttrs.addAttribute(llvm::Attribute::NoUnwind);
+
     // Respect -fcuda-flush-denormals-to-zero.
     if (getLangOpts().CUDADeviceFlushDenormalsToZero)
       FuncAttrs.addAttribute("nvptx-f32ftz", "true");
diff --git a/clang/lib/CodeGen/CGException.cpp b/clang/lib/CodeGen/CGException.cpp
index 4a7dc4205e09..57ac44e6b174 100644
--- a/clang/lib/CodeGen/CGException.cpp
+++ b/clang/lib/CodeGen/CGException.cpp
@@ -698,6 +698,10 @@ llvm::BasicBlock *CodeGenFunction::getInvokeDestImpl() {
       return nullptr;
   }
 
+  // CUDA device code doesn't have exceptions.
+  if (LO.CUDA && LO.CUDAIsDevice)
+    return nullptr;
+
   // Check the innermost scope for a cached landing pad.  If this is
   // a non-EH cleanup, we'll check enclosing scopes in EmitLandingPad.
   llvm::BasicBlock *LP = EHStack.begin()->getCachedLandingPad();
diff --git a/clang/test/CodeGenCUDA/convergent.cu b/clang/test/CodeGenCUDA/convergent.cu
index 6827c57d29fb..62818f9e5af4 100644
--- a/clang/test/CodeGenCUDA/convergent.cu
+++ b/clang/test/CodeGenCUDA/convergent.cu
@@ -36,8 +36,8 @@ __host__ __device__ void bar() {
 // DEVICE: attributes [[BAZ_ATTR]] = {
 // DEVICE-SAME: convergent
 // DEVICE-SAME: }
-// DEVICE: attributes [[CALL_ATTR]] = { convergent }
-// DEVICE: attributes [[ASM_ATTR]] = { convergent
+// DEVICE-DAG: attributes [[CALL_ATTR]] = { convergent
+// DEVICE-DAG: attributes [[ASM_ATTR]] = { convergent
 
 // HOST: declare void @_Z3bazv() [[BAZ_ATTR:#[0-9]+]]
 // HOST: attributes [[BAZ_ATTR]] = {
diff --git a/clang/test/CodeGenCUDA/device-var-init.cu b/clang/test/CodeGenCUDA/device-var-init.cu
index 6f2d9294131f..0f4c64813406 100644
--- a/clang/test/CodeGenCUDA/device-var-init.cu
+++ b/clang/test/CodeGenCUDA/device-var-init.cu
@@ -182,9 +182,9 @@ __device__ void df() {
   df(); // CHECK: call void @_Z2dfv()
 
   // Verify that we only call non-empty destructors
-  // CHECK-NEXT: call void @_ZN8T_FA_NEDD1Ev(%struct.T_FA_NED* %t_fa_ned) #6
-  // CHECK-NEXT: call void @_ZN7T_F_NEDD1Ev(%struct.T_F_NED* %t_f_ned) #6
-  // CHECK-NEXT: call void @_ZN7T_B_NEDD1Ev(%struct.T_B_NED* %t_b_ned) #6
+  // CHECK-NEXT: call void @_ZN8T_FA_NEDD1Ev(%struct.T_FA_NED* %t_fa_ned)
+  // CHECK-NEXT: call void @_ZN7T_F_NEDD1Ev(%struct.T_F_NED* %t_f_ned)
+  // CHECK-NEXT: call void @_ZN7T_B_NEDD1Ev(%struct.T_B_NED* %t_b_ned)
   // CHECK-NEXT: call void @_ZN2VDD1Ev(%struct.VD* %vd)
   // CHECK-NEXT: call void @_ZN3NEDD1Ev(%struct.NED* %ned)
   // CHECK-NEXT: call void @_ZN2UDD1Ev(%struct.UD* %ud)
diff --git a/clang/test/CodeGenCUDA/nothrow.cu b/clang/test/CodeGenCUDA/nothrow.cu
new file mode 100644
index 000000000000..aa5bfcdda6a8
--- /dev/null
+++ b/clang/test/CodeGenCUDA/nothrow.cu
@@ -0,0 +1,39 @@
+// RUN: %clang_cc1 -std=c++11 -fcxx-exceptions -fexceptions -fcuda-is-device \
+// RUN:   -triple nvptx-nvidia-cuda -emit-llvm -disable-llvm-passes -o - %s | \
+// RUN  FileCheck -check-prefix DEVICE %s
+
+// RUN: %clang_cc1 -std=c++11 -fcxx-exceptions -fexceptions \
+// RUN:   -triple x86_64-unknown-linux-gnu -emit-llvm -disable-llvm-passes -o - %s | \
+// RUN:  FileCheck -check-prefix HOST %s
+
+#include "Inputs/cuda.h"
+
+__host__ __device__ void f();
+
+// HOST: define void @_Z7host_fnv() [[HOST_ATTR:#[0-9]+]]
+void host_fn() { f(); }
+
+// DEVICE: define void @_Z3foov() [[DEVICE_ATTR:#[0-9]+]]
+__device__ void foo() {
+  // DEVICE: call void @_Z1fv
+  f();
+}
+
+// DEVICE: define void @_Z12foo_noexceptv() [[DEVICE_ATTR:#[0-9]+]]
+__device__ void foo_noexcept() noexcept {
+  // DEVICE: call void @_Z1fv
+  f();
+}
+
+// This is nounwind only on the device side.
+// CHECK: define void @_Z3foov() [[DEVICE_ATTR:#[0-9]+]]
+__host__ __device__ void bar() { f(); }
+
+// DEVICE: define void @_Z3bazv() [[DEVICE_ATTR:#[0-9]+]]
+__global__ void baz() { f(); }
+
+// DEVICE: attributes [[DEVICE_ATTR]] = {
+// DEVICE-SAME: nounwind
+// HOST: attributes [[HOST_ATTR]] = {
+// HOST-NOT: nounwind
+// HOST-SAME: }