diff --git a/clang/include/clang/Basic/DiagnosticCommonKinds.td b/clang/include/clang/Basic/DiagnosticCommonKinds.td
index cd72e254ea3b1..f2df283c74829 100644
--- a/clang/include/clang/Basic/DiagnosticCommonKinds.td
+++ b/clang/include/clang/Basic/DiagnosticCommonKinds.td
@@ -425,4 +425,13 @@ def err_opencl_extension_and_feature_differs : Error<
   "options %0 and %1 are set to different values">;
 def err_opencl_feature_requires : Error<
   "feature %0 requires support of %1 feature">;
+
+def warn_throw_not_valid_on_target : Warning<
+  "target '%0' does not support exception handling;"
+  " 'throw' is assumed to be never reached">,
+  InGroup<OpenMPTargetException>;
+def warn_try_not_valid_on_target : Warning<
+  "target '%0' does not support exception handling;"
+  " 'catch' block is ignored">,
+  InGroup<OpenMPTargetException>;
 }
diff --git a/clang/include/clang/Basic/DiagnosticGroups.td b/clang/include/clang/Basic/DiagnosticGroups.td
index d1aa51393ef35..00c458fb23e73 100644
--- a/clang/include/clang/Basic/DiagnosticGroups.td
+++ b/clang/include/clang/Basic/DiagnosticGroups.td
@@ -1292,9 +1292,10 @@ def OpenMPTarget : DiagGroup<"openmp-target", [OpenMPMapping]>;
 def OpenMPPre51Compat : DiagGroup<"pre-openmp-51-compat">;
 def OpenMP51Ext : DiagGroup<"openmp-51-extensions">;
 def OpenMPExtensions : DiagGroup<"openmp-extensions">;
+def OpenMPTargetException : DiagGroup<"openmp-target-exception">;
 def OpenMP : DiagGroup<"openmp", [
     SourceUsesOpenMP, OpenMPClauses, OpenMPLoopForm, OpenMPTarget,
-    OpenMPMapping, OpenMP51Ext, OpenMPExtensions
+    OpenMPMapping, OpenMP51Ext, OpenMPExtensions, OpenMPTargetException
   ]>;
 
 // Backend warnings.
diff --git a/clang/lib/CodeGen/CGException.cpp b/clang/lib/CodeGen/CGException.cpp
index 9cb7d4c7731de..3996f2948349c 100644
--- a/clang/lib/CodeGen/CGException.cpp
+++ b/clang/lib/CodeGen/CGException.cpp
@@ -440,6 +440,15 @@ llvm::Value *CodeGenFunction::getSelectorFromSlot() {
 
 void CodeGenFunction::EmitCXXThrowExpr(const CXXThrowExpr *E,
                                        bool KeepInsertionPoint) {
+  // If the exception is being emitted in an OpenMP target region,
+  // and the target is a GPU, we do not support exception handling.
+  // Therefore, we emit a trap which will abort the program, and
+  // prompt a warning indicating that a trap will be emitted.
+  const llvm::Triple &T = Target.getTriple();
+  if (CGM.getLangOpts().OpenMPIsTargetDevice && (T.isNVPTX() || T.isAMDGCN())) {
+    EmitTrapCall(llvm::Intrinsic::trap);
+    return;
+  }
   if (const Expr *SubExpr = E->getSubExpr()) {
     QualType ThrowType = SubExpr->getType();
     if (ThrowType->isObjCObjectPointerType()) {
@@ -609,9 +618,16 @@ void CodeGenFunction::EmitEndEHSpec(const Decl *D) {
 }
 
 void CodeGenFunction::EmitCXXTryStmt(const CXXTryStmt &S) {
-  EnterCXXTryStmt(S);
+  const llvm::Triple &T = Target.getTriple();
+  // If we encounter a try statement on in an OpenMP target region offloaded to
+  // a GPU, we treat it as a basic block.
+  const bool IsTargetDevice =
+      (CGM.getLangOpts().OpenMPIsTargetDevice && (T.isNVPTX() || T.isAMDGCN()));
+  if (!IsTargetDevice)
+    EnterCXXTryStmt(S);
   EmitStmt(S.getTryBlock());
-  ExitCXXTryStmt(S);
+  if (!IsTargetDevice)
+    ExitCXXTryStmt(S);
 }
 
 void CodeGenFunction::EnterCXXTryStmt(const CXXTryStmt &S, bool IsFnTryBlock) {
diff --git a/clang/lib/Sema/SemaExprCXX.cpp b/clang/lib/Sema/SemaExprCXX.cpp
index a6e0a8abf81b9..472fbdbdb5d0e 100644
--- a/clang/lib/Sema/SemaExprCXX.cpp
+++ b/clang/lib/Sema/SemaExprCXX.cpp
@@ -864,13 +864,21 @@ Sema::ActOnCXXThrow(Scope *S, SourceLocation OpLoc, Expr *Ex) {
 
 ExprResult Sema::BuildCXXThrow(SourceLocation OpLoc, Expr *Ex,
                                bool IsThrownVarInScope) {
-  // Don't report an error if 'throw' is used in system headers.
-  if (!getLangOpts().CXXExceptions &&
+  const llvm::Triple &T = Context.getTargetInfo().getTriple();
+  const bool IsOpenMPGPUTarget =
+      getLangOpts().OpenMPIsTargetDevice && (T.isNVPTX() || T.isAMDGCN());
+  // Don't report an error if 'throw' is used in system headers or in an OpenMP
+  // target region compiled for a GPU architecture.
+  if (!IsOpenMPGPUTarget && !getLangOpts().CXXExceptions &&
       !getSourceManager().isInSystemHeader(OpLoc) && !getLangOpts().CUDA) {
     // Delay error emission for the OpenMP device code.
     targetDiag(OpLoc, diag::err_exceptions_disabled) << "throw";
   }
 
+  // In OpenMP target regions, we replace 'throw' with a trap on GPU targets.
+  if (IsOpenMPGPUTarget)
+    targetDiag(OpLoc, diag::warn_throw_not_valid_on_target) << T.str();
+
   // Exceptions aren't allowed in CUDA device code.
   if (getLangOpts().CUDA)
     CUDADiagIfDeviceCode(OpLoc, diag::err_cuda_device_exceptions)
diff --git a/clang/lib/Sema/SemaStmt.cpp b/clang/lib/Sema/SemaStmt.cpp
index 70a549938d080..88ea3f1c3349d 100644
--- a/clang/lib/Sema/SemaStmt.cpp
+++ b/clang/lib/Sema/SemaStmt.cpp
@@ -4471,13 +4471,22 @@ class CatchTypePublicBases {
 /// handlers and creates a try statement from them.
 StmtResult Sema::ActOnCXXTryBlock(SourceLocation TryLoc, Stmt *TryBlock,
                                   ArrayRef<Stmt *> Handlers) {
-  // Don't report an error if 'try' is used in system headers.
-  if (!getLangOpts().CXXExceptions &&
+  const llvm::Triple &T = Context.getTargetInfo().getTriple();
+  const bool IsOpenMPGPUTarget =
+      getLangOpts().OpenMPIsTargetDevice && (T.isNVPTX() || T.isAMDGCN());
+  // Don't report an error if 'try' is used in system headers or in an OpenMP
+  // target region compiled for a GPU architecture.
+  if (!IsOpenMPGPUTarget && !getLangOpts().CXXExceptions &&
       !getSourceManager().isInSystemHeader(TryLoc) && !getLangOpts().CUDA) {
     // Delay error emission for the OpenMP device code.
     targetDiag(TryLoc, diag::err_exceptions_disabled) << "try";
   }
 
+  // In OpenMP target regions, we assume that catch is never reached on GPU
+  // targets.
+  if (IsOpenMPGPUTarget)
+    targetDiag(TryLoc, diag::warn_try_not_valid_on_target) << T.str();
+
   // Exceptions aren't allowed in CUDA device code.
   if (getLangOpts().CUDA)
     CUDADiagIfDeviceCode(TryLoc, diag::err_cuda_device_exceptions)
diff --git a/clang/test/OpenMP/amdgpu_exceptions.cpp b/clang/test/OpenMP/amdgpu_exceptions.cpp
new file mode 100644
index 0000000000000..f381ec30b45a0
--- /dev/null
+++ b/clang/test/OpenMP/amdgpu_exceptions.cpp
@@ -0,0 +1,46 @@
+/**
+ * The first four lines test that a warning is produced when enabling 
+ * -Wopenmp-target-exception no matter what combination of -fexceptions and 
+ * -fcxx-exceptions are set, as we want OpenMP to always allow exceptions in the
+ * target region but emit a warning instead.
+*/
+
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+
+/**
+ * The following four lines test that no warning is emitted when providing 
+ * -Wno-openmp-target-exception no matter the combination of -fexceptions and 
+ * -fcxx-exceptions.
+*/
+
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fcxx-exceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fexceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+
+/**
+ * Finally we should test that we only ignore exceptions in the OpenMP 
+ * offloading tool-chain
+*/
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa %s -emit-llvm -S -verify=noexceptions -o -
+
+// noexceptions-error@37 {{cannot use 'try' with exceptions disabled}}
+// noexceptions-error@38 {{cannot use 'throw' with exceptions disabled}}
+
+#pragma omp declare target
+int foo(void) {
+	int error = -1;
+	try { // with-warning {{target 'amdgcn-amd-amdhsa' does not support exception handling; 'catch' block is ignored}}
+		throw 404; // with-warning {{target 'amdgcn-amd-amdhsa' does not support exception handling; 'throw' is assumed to be never reached}}
+	}
+	catch (int e){ 
+		error = e;
+	}
+	return error;
+}
+#pragma omp end declare target
+// without-no-diagnostics
diff --git a/clang/test/OpenMP/amdgpu_throw.cpp b/clang/test/OpenMP/amdgpu_throw.cpp
new file mode 100644
index 0000000000000..bccba270d05bc
--- /dev/null
+++ b/clang/test/OpenMP/amdgpu_throw.cpp
@@ -0,0 +1,38 @@
+/**
+ * The first four lines test that a warning is produced when enabling 
+ * -Wopenmp-target-exception no matter what combination of -fexceptions and 
+ * -fcxx-exceptions are set, as we want OpenMP to always allow exceptions in the
+ * target region but emit a warning instead.
+*/
+
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+
+/**
+ * The following four lines test that no warning is emitted when providing 
+ * -Wno-openmp-target-exception no matter the combination of -fexceptions and 
+ * -fcxx-exceptions.
+*/
+
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fcxx-exceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fexceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+
+/**
+ * Finally we should test that we only ignore exceptions in the OpenMP 
+ * offloading tool-chain
+*/
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa %s -emit-llvm -S -verify=noexceptions -o -
+
+// noexceptions-error@35 {{cannot use 'throw' with exceptions disabled}}
+
+#pragma omp declare target
+void foo(void) {
+	throw 404; // with-warning {{target 'amdgcn-amd-amdhsa' does not support exception handling; 'throw' is assumed to be never reached}}
+}
+#pragma omp end declare target
+// without-no-diagnostics
diff --git a/clang/test/OpenMP/amdgpu_throw_trap.cpp b/clang/test/OpenMP/amdgpu_throw_trap.cpp
new file mode 100644
index 0000000000000..513b543ccaa8f
--- /dev/null
+++ b/clang/test/OpenMP/amdgpu_throw_trap.cpp
@@ -0,0 +1,11 @@
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device %s -emit-llvm -S -Wno-openmp-target-exception -o - | FileCheck -check-prefix=DEVICE %s
+// RUN: %clang_cc1 -fopenmp -triple x86_64-pc-linux-gnu -fopenmp-is-target-device -fcxx-exceptions %s -emit-llvm -S -Wno-openmp-target-exception -o - | FileCheck -check-prefix=HOST %s
+// DEVICE: s_trap
+// DEVICE-NOT: __cxa_throw
+// HOST: __cxa_throw
+// HOST-NOT: s_trap
+#pragma omp declare target
+void foo(void) {
+	throw 404; 
+}
+#pragma omp end declare target
diff --git a/clang/test/OpenMP/amdgpu_try_catch.cpp b/clang/test/OpenMP/amdgpu_try_catch.cpp
new file mode 100644
index 0000000000000..7aa20419edcd4
--- /dev/null
+++ b/clang/test/OpenMP/amdgpu_try_catch.cpp
@@ -0,0 +1,45 @@
+/**
+ * The first four lines test that a warning is produced when enabling 
+ * -Wopenmp-target-exception no matter what combination of -fexceptions and 
+ * -fcxx-exceptions are set, as we want OpenMP to always allow exceptions in the
+ * target region but emit a warning instead.
+*/
+
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+
+/**
+ * The following four lines test that no warning is emitted when providing 
+ * -Wno-openmp-target-exception no matter the combination of -fexceptions and 
+ * -fcxx-exceptions.
+*/
+
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fcxx-exceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device -fexceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple amdgcn-amd-amdhsa -fopenmp-is-target-device %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+
+/**
+ * Finally we should test that we only ignore exceptions in the OpenMP 
+ * offloading tool-chain
+*/
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa %s -emit-llvm -S -verify=noexceptions -o -
+
+// noexceptions-error@36 {{cannot use 'try' with exceptions disabled}}
+
+#pragma omp declare target
+int foo(void) {
+	int error = -1;
+	try { // with-warning {{target 'amdgcn-amd-amdhsa' does not support exception handling; 'catch' block is ignored}}
+		error = 1;
+	}
+	catch (int e){ 
+		error = e;
+	}
+	return error;
+}
+#pragma omp end declare target
+// without-no-diagnostics
diff --git a/clang/test/OpenMP/nvptx_exceptions.cpp b/clang/test/OpenMP/nvptx_exceptions.cpp
new file mode 100644
index 0000000000000..e59fb1f3d30bc
--- /dev/null
+++ b/clang/test/OpenMP/nvptx_exceptions.cpp
@@ -0,0 +1,46 @@
+/**
+ * The first four lines test that a warning is produced when enabling 
+ * -Wopenmp-target-exception no matter what combination of -fexceptions and 
+ * -fcxx-exceptions are set, as we want OpenMP to always allow exceptions in the
+ * target region but emit a warning instead.
+*/
+
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+
+/**
+ * The following four lines test that no warning is emitted when providing 
+ * -Wno-openmp-target-exception no matter the combination of -fexceptions and 
+ * -fcxx-exceptions.
+*/
+
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fcxx-exceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fexceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+
+/**
+ * Finally we should test that we only ignore exceptions in the OpenMP 
+ * offloading tool-chain
+*/
+
+// RUN: %clang_cc1 -triple nvptx64 %s -emit-llvm -S -verify=noexceptions -o -
+
+// noexceptions-error@37 {{cannot use 'try' with exceptions disabled}}
+// noexceptions-error@38 {{cannot use 'throw' with exceptions disabled}}
+
+#pragma omp declare target
+int foo(void) {
+	int error = -1;
+	try { // with-warning {{target 'nvptx64' does not support exception handling; 'catch' block is ignored}}
+		throw 404; // with-warning {{target 'nvptx64' does not support exception handling; 'throw' is assumed to be never reached}}
+	}
+	catch (int e){ 
+		error = e;
+	}
+	return error;
+}
+#pragma omp end declare target
+// without-no-diagnostics
diff --git a/clang/test/OpenMP/nvptx_target_exceptions_messages.cpp b/clang/test/OpenMP/nvptx_target_exceptions_messages.cpp
index 9f267f5606e2d..5d1d46cadff26 100644
--- a/clang/test/OpenMP/nvptx_target_exceptions_messages.cpp
+++ b/clang/test/OpenMP/nvptx_target_exceptions_messages.cpp
@@ -34,7 +34,7 @@ T FA() {
 #pragma omp declare target
 struct S {
   int a;
-  S(int a) : a(a) { throw 1; } // expected-error {{cannot use 'throw' with exceptions disabled}}
+  S(int a) : a(a) { throw 1; } // expected-warning {{target 'nvptx64-unknown-unknown' does not support exception handling; 'throw' is assumed to be never reached}}
 };
 
 int foo() { return 0; }
@@ -57,7 +57,7 @@ int maini1() {
     static long aaa = 23;
     a = foo() + bar() + b + c + d + aa + aaa + FA<int>(); // expected-note{{called by 'maini1'}}
     if (!a)
-      throw "Error"; // expected-error {{cannot use 'throw' with exceptions disabled}}
+      throw "Error"; // expected-warning {{target 'nvptx64-unknown-unknown' does not support exception handling; 'throw' is assumed to be never reached}}
   }
   } catch(...) {
   }
@@ -67,14 +67,14 @@ int maini1() {
 int baz3() { return 2 + baz2(); }
 int baz2() {
 #pragma omp target
-  try { // expected-error {{cannot use 'try' with exceptions disabled}}
+  try { // expected-warning {{target 'nvptx64-unknown-unknown' does not support exception handling; 'catch' block is ignored}}
   ++c;
   } catch (...) {
   }
   return 2 + baz3();
 }
 
-int baz1() { throw 1; } // expected-error {{cannot use 'throw' with exceptions disabled}}
+int baz1() { throw 1; } // expected-warning {{target 'nvptx64-unknown-unknown' does not support exception handling; 'throw' is assumed to be never reached}}
 
 int foobar1();
 int foobar2();
@@ -85,7 +85,7 @@ int (*B)() = &foobar2;
 #pragma omp end declare target
 
 int foobar1() { throw 1; }
-int foobar2() { throw 1; } // expected-error {{cannot use 'throw' with exceptions disabled}}
+int foobar2() { throw 1; } // expected-warning {{target 'nvptx64-unknown-unknown' does not support exception handling; 'throw' is assumed to be never reached}}
 
 
 int foobar3();
@@ -95,7 +95,7 @@ int (*C)() = &foobar3; // expected-warning {{declaration is not declared in any
 int (*D)() = C; // expected-note {{used here}}
                 // host-note@-1 {{used here}}
 #pragma omp end declare target
-int foobar3() { throw 1; } // expected-error {{cannot use 'throw' with exceptions disabled}}
+int foobar3() { throw 1; } // expected-warning {{target 'nvptx64-unknown-unknown' does not support exception handling; 'throw' is assumed to be never reached}}
 
 // Check no infinite recursion in deferred diagnostic emitter.
 long E = (long)&E;
diff --git a/clang/test/OpenMP/nvptx_throw.cpp b/clang/test/OpenMP/nvptx_throw.cpp
new file mode 100644
index 0000000000000..6d5774551fe94
--- /dev/null
+++ b/clang/test/OpenMP/nvptx_throw.cpp
@@ -0,0 +1,38 @@
+/**
+ * The first four lines test that a warning is produced when enabling 
+ * -Wopenmp-target-exception no matter what combination of -fexceptions and 
+ * -fcxx-exceptions are set, as we want OpenMP to always allow exceptions in the
+ * target region but emit a warning instead.
+*/
+
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+
+/**
+ * The following four lines test that no warning is emitted when providing 
+ * -Wno-openmp-target-exception no matter the combination of -fexceptions and 
+ * -fcxx-exceptions.
+*/
+
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fcxx-exceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fexceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+
+/**
+ * Finally we should test that we only ignore exceptions in the OpenMP 
+ * offloading tool-chain
+*/
+
+// RUN: %clang_cc1 -triple nvptx64 %s -emit-llvm -S -verify=noexceptions -o -
+
+// noexceptions-error@35 {{cannot use 'throw' with exceptions disabled}}
+
+#pragma omp declare target
+void foo(void) {
+	throw 404; // with-warning {{target 'nvptx64' does not support exception handling; 'throw' is assumed to be never reached}}
+}
+#pragma omp end declare target
+// without-no-diagnostics
diff --git a/clang/test/OpenMP/nvptx_throw_trap.cpp b/clang/test/OpenMP/nvptx_throw_trap.cpp
new file mode 100644
index 0000000000000..d09f4dfd88571
--- /dev/null
+++ b/clang/test/OpenMP/nvptx_throw_trap.cpp
@@ -0,0 +1,11 @@
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device %s -emit-llvm -S -Wno-openmp-target-exception -o - | FileCheck -check-prefix=DEVICE %s
+// RUN: %clang_cc1 -fopenmp -triple x86_64-pc-linux-gnu -fopenmp-is-target-device -fcxx-exceptions %s -emit-llvm -S -Wno-openmp-target-exception -o - | FileCheck -check-prefix=HOST %s
+// DEVICE: trap;
+// DEVICE-NOT: __cxa_throw
+// HOST: __cxa_throw
+// HOST-NOT: trap;
+#pragma omp declare target
+void foo(void) {
+	throw 404; 
+}
+#pragma omp end declare target
diff --git a/clang/test/OpenMP/nvptx_try_catch.cpp b/clang/test/OpenMP/nvptx_try_catch.cpp
new file mode 100644
index 0000000000000..c2a0786aefd53
--- /dev/null
+++ b/clang/test/OpenMP/nvptx_try_catch.cpp
@@ -0,0 +1,45 @@
+/**
+ * The first four lines test that a warning is produced when enabling 
+ * -Wopenmp-target-exception no matter what combination of -fexceptions and 
+ * -fcxx-exceptions are set, as we want OpenMP to always allow exceptions in the
+ * target region but emit a warning instead.
+*/
+
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fexceptions %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device %s -emit-llvm -S -verify=with -Wopenmp-target-exception -analyze
+
+/**
+ * The following four lines test that no warning is emitted when providing 
+ * -Wno-openmp-target-exception no matter the combination of -fexceptions and 
+ * -fcxx-exceptions.
+*/
+
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fcxx-exceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device -fexceptions %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+// RUN: %clang_cc1 -fopenmp -triple nvptx64 -fopenmp-is-target-device %s -emit-llvm -S -verify=without -Wno-openmp-target-exception -analyze
+
+/**
+ * Finally we should test that we only ignore exceptions in the OpenMP 
+ * offloading tool-chain
+*/
+
+// RUN: %clang_cc1 -triple nvptx64 %s -emit-llvm -S -verify=noexceptions -o -
+
+// noexceptions-error@36 {{cannot use 'try' with exceptions disabled}}
+
+#pragma omp declare target
+int foo(void) {
+	int error = -1;
+	try { // with-warning {{target 'nvptx64' does not support exception handling; 'catch' block is ignored}}
+		error = 1;
+	}
+	catch (int e){ 
+		error = e;
+	}
+	return error;
+}
+#pragma omp end declare target
+// without-no-diagnostics
diff --git a/clang/test/OpenMP/x86_target_exceptions.cpp b/clang/test/OpenMP/x86_target_exceptions.cpp
new file mode 100644
index 0000000000000..bc5e485b88345
--- /dev/null
+++ b/clang/test/OpenMP/x86_target_exceptions.cpp
@@ -0,0 +1,14 @@
+// RUN: %clang_cc1 -fopenmp -triple x86_64-pc-linux-gnu -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify -Wopenmp-target-exception -analyze
+#pragma omp declare target
+int foo(void) {
+	int error = -1;
+	try {
+		throw 404;
+	}
+	catch (int e){ 
+		error = e;
+	}
+	return error;
+}
+#pragma omp end declare target
+// expected-no-diagnostics
diff --git a/clang/test/OpenMP/x86_target_throw.cpp b/clang/test/OpenMP/x86_target_throw.cpp
new file mode 100644
index 0000000000000..03fef3faa83d9
--- /dev/null
+++ b/clang/test/OpenMP/x86_target_throw.cpp
@@ -0,0 +1,7 @@
+// RUN: %clang_cc1 -fopenmp -triple x86_64-pc-linux-gnu -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify -Wopenmp-target-exception -analyze
+#pragma omp declare target
+void foo(void) {
+	throw 404;
+}
+#pragma omp end declare target
+// expected-no-diagnostics
diff --git a/clang/test/OpenMP/x86_target_try_catch.cpp b/clang/test/OpenMP/x86_target_try_catch.cpp
new file mode 100644
index 0000000000000..e878d2ee0929b
--- /dev/null
+++ b/clang/test/OpenMP/x86_target_try_catch.cpp
@@ -0,0 +1,14 @@
+// RUN: %clang_cc1 -fopenmp -triple x86_64-pc-linux-gnu -fopenmp-is-target-device -fcxx-exceptions -fexceptions %s -emit-llvm -S -verify -Wopenmp-target-exception -analyze
+#pragma omp declare target
+int foo(void) {
+	int error = -1;
+	try {
+		error = 1;
+	}
+	catch (int e){ 
+		error = e;
+	}
+	return error;
+}
+#pragma omp end declare target
+// expected-no-diagnostics