diff --git a/clang/include/clang/Basic/DiagnosticCommonKinds.td b/clang/include/clang/Basic/DiagnosticCommonKinds.td
--- a/clang/include/clang/Basic/DiagnosticCommonKinds.td
+++ b/clang/include/clang/Basic/DiagnosticCommonKinds.td
@@ -420,4 +420,14 @@
   "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."
+  " To allow code generation for '%0', 'throw' expressions will be replaced by traps.">;
+def warn_try_not_valid_on_target : Warning<
+  "Target '%0' does not support exception handling."
+  " To allow code generation for '%0', 'try' statements are treated as basic blocks.">;
+def warn_catch_not_valid_on_target : Warning<
+  "Target '%0' does not support exception handling."
+  " To allow code generation for '%0', 'catch' statement will be replaced by a no operation instruction.">;
 }
diff --git a/clang/lib/CodeGen/CGException.cpp b/clang/lib/CodeGen/CGException.cpp
--- a/clang/lib/CodeGen/CGException.cpp
+++ b/clang/lib/CodeGen/CGException.cpp
@@ -9,7 +9,6 @@
 // This contains code dealing with C++ exception related code generation.
 //
 //===----------------------------------------------------------------------===//
-
 #include "CGCXXABI.h"
 #include "CGCleanup.h"
 #include "CGObjCRuntime.h"
@@ -440,7 +439,16 @@
 
 void CodeGenFunction::EmitCXXThrowExpr(const CXXThrowExpr *E,
                                        bool KeepInsertionPoint) {
-  if (const Expr *SubExpr = E->getSubExpr()) {
+  // 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().OpenMPIsDevice && (T.isNVPTX() || T.isAMDGCN())) {
+    CGM.getDiags().Report(E->getExprLoc(), diag::warn_throw_not_valid_on_target)
+        << T.str();
+    EmitTrapCall(llvm::Intrinsic::trap);
+  } else if (const Expr *SubExpr = E->getSubExpr()) {
     QualType ThrowType = SubExpr->getType();
     if (ThrowType->isObjCObjectPointerType()) {
       const Stmt *ThrowStmt = E->getSubExpr();
@@ -609,9 +617,20 @@
 }
 
 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 is_omp_gpu_target =
+      (CGM.getLangOpts().OpenMPIsDevice && (T.isNVPTX() || T.isAMDGCN()));
+  if (is_omp_gpu_target) {
+    CGM.getDiags().Report(S.getTryLoc(), diag::warn_try_not_valid_on_target)
+        << T.str();
+  }
+  if (!is_omp_gpu_target)
+    EnterCXXTryStmt(S);
   EmitStmt(S.getTryBlock());
-  ExitCXXTryStmt(S);
+  if (!is_omp_gpu_target)
+    ExitCXXTryStmt(S);
 }
 
 void CodeGenFunction::EnterCXXTryStmt(const CXXTryStmt &S, bool IsFnTryBlock) {
diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp
--- a/clang/lib/Frontend/CompilerInvocation.cpp
+++ b/clang/lib/Frontend/CompilerInvocation.cpp
@@ -3826,12 +3826,15 @@
   }
 
   // Set the flag to prevent the implementation from emitting device exception
-  // handling code for those requiring so.
+  // handling code for those requiring so. However, if the user explicitly
+  // enabled exception handling on the device, we will allow exceptions during
+  // Sema and handle the exceptions differently in CodeGen.
   if ((Opts.OpenMPIsDevice && (T.isNVPTX() || T.isAMDGCN())) ||
       Opts.OpenCLCPlusPlus) {
-
-    Opts.Exceptions = 0;
-    Opts.CXXExceptions = 0;
+    bool exceptions_user_enabled = Args.hasFlag(
+        options::OPT_fexceptions, options::OPT_fno_exceptions, false);
+    Opts.Exceptions = exceptions_user_enabled;
+    Opts.CXXExceptions = exceptions_user_enabled;
   }
   if (Opts.OpenMPIsDevice && T.isNVPTX()) {
     Opts.OpenMPCUDANumSMs =
diff --git a/clang/test/OpenMP/target_throw_message.cpp b/clang/test/OpenMP/target_throw_message.cpp
new file mode 100644
--- /dev/null
+++ b/clang/test/OpenMP/target_throw_message.cpp
@@ -0,0 +1,27 @@
+// We first test that we treat 'throw' as a trap on AMD and NVIDIA GPUs in TEST1 and TEST2
+// RUN: %clangxx -O0 -Wall -std=c++17 -fopenmp \
+// RUN: -fopenmp-targets=amdgcn-amd-amdhsa --offload-arch=gfx906 \
+// RUN: -fopenmp-offload-mandatory -Xopenmp-target -fexceptions \
+// RUN: -Xopenmp-target -fcxx-exceptions -emit-llvm -S %s -o - 2>&1 | FileCheck %s -v --check-prefix=TEST1
+// TEST1: {{Target 'amdgcn-amd-amdhsa' does not support exception handling. To allow code generation for 'amdgcn-amd-amdhsa', 'throw' expressions will be replaced by traps}}
+
+// RUN: %clangxx -O0 -Wall -std=c++17 -fopenmp \
+// RUN: -fopenmp-targets=nvptx64-nvidia-cuda --offload-arch=sm_70 \
+// RUN: -fopenmp-offload-mandatory -Xopenmp-target -fexceptions \
+// RUN: -Xopenmp-target -fcxx-exceptions -emit-llvm -S %s -o - 2>&1 | FileCheck %s -v --check-prefix=TEST2
+// TEST2: {{Target 'nvptx64-nvidia-cuda' does not support exception handling. To allow code generation for 'nvptx64-nvidia-cuda', 'throw' expressions will be replaced by traps}}
+
+// Then, we test that exception handling is still allowed on code offloaded to a CPU in TEST3
+// RUN: %clangxx -O0 -Wall -std=c++17 -fopenmp \
+// RUN: -fopenmp-targets=x86_64-pc-linux-gnu --offload-arch=x86_64-pc-linux-gnu \
+// RUN: -fopenmp-offload-mandatory -Xopenmp-target -fexceptions \
+// RUN: -Xopenmp-target -fcxx-exceptions -emit-llvm -S %s -o - 2>&1 | FileCheck %s -v --check-prefix=TEST3
+// TEST3-NOT: {{Target 'x86_64-pc-linux-gnu' does not support exception handling. To allow code generation for 'x86_64-pc-linux-gnu', 'throw' expressions will be replaced by traps}}
+
+int main(void) {
+#pragma omp target
+  {
+    throw 404;
+  }
+  return 0;
+}
\ No newline at end of file
diff --git a/clang/test/OpenMP/target_throw_message_fun_call.cpp b/clang/test/OpenMP/target_throw_message_fun_call.cpp
new file mode 100644
--- /dev/null
+++ b/clang/test/OpenMP/target_throw_message_fun_call.cpp
@@ -0,0 +1,32 @@
+// We first test that we treat 'throw' as a trap on AMD and NVIDIA GPUs in TEST1 and TEST2
+// RUN: %clangxx -O0 -Wall -std=c++17 -fopenmp \
+// RUN: -fopenmp-targets=amdgcn-amd-amdhsa --offload-arch=gfx906 \
+// RUN: -fopenmp-offload-mandatory -Xopenmp-target -fexceptions \
+// RUN: -Xopenmp-target -fcxx-exceptions -emit-llvm -S %s -o - 2>&1 | FileCheck %s -v --check-prefix=TEST1
+// TEST1: {{Target 'amdgcn-amd-amdhsa' does not support exception handling. To allow code generation for 'amdgcn-amd-amdhsa', 'throw' expressions will be replaced by traps}}
+
+// RUN: %clangxx -O0 -Wall -std=c++17 -fopenmp \
+// RUN: -fopenmp-targets=nvptx64-nvidia-cuda --offload-arch=sm_70 \
+// RUN: -fopenmp-offload-mandatory -Xopenmp-target -fexceptions \
+// RUN: -Xopenmp-target -fcxx-exceptions -emit-llvm -S %s -o - 2>&1 | FileCheck %s -v --check-prefix=TEST2
+// TEST2: {{Target 'nvptx64-nvidia-cuda' does not support exception handling. To allow code generation for 'nvptx64-nvidia-cuda', 'throw' expressions will be replaced by traps}}
+
+// Then, we test that exception handling is still allowed on code offloaded to a CPU in TEST3
+// RUN: %clangxx -O0 -Wall -std=c++17 -fopenmp \
+// RUN: -fopenmp-targets=x86_64-pc-linux-gnu --offload-arch=x86_64-pc-linux-gnu \
+// RUN: -fopenmp-offload-mandatory -Xopenmp-target -fexceptions \
+// RUN: -Xopenmp-target -fcxx-exceptions -emit-llvm -S %s -o - 2>&1 | FileCheck %s -v --check-prefix=TEST3
+// TEST3-NOT: {{Target 'x86_64-pc-linux-gnu' does not support exception handling. To allow code generation for 'x86_64-pc-linux-gnu', 'throw' expressions will be replaced by traps}}
+
+
+#pragma omp declare target
+void foo() {throw 'e';}
+#pragma omp end declare target
+
+int main(void) {
+#pragma omp target
+  {
+    foo();
+  }
+  return 0;
+}
\ No newline at end of file
diff --git a/clang/test/OpenMP/target_try_catch_message.cpp b/clang/test/OpenMP/target_try_catch_message.cpp
new file mode 100644
--- /dev/null
+++ b/clang/test/OpenMP/target_try_catch_message.cpp
@@ -0,0 +1,33 @@
+// We first test that we treat 'try' as a basic block on AMD and NVIDIA GPUs in TEST1 and TEST2
+// RUN: %clangxx -O0 -Wall -std=c++17 -fopenmp \
+// RUN: -fopenmp-targets=amdgcn-amd-amdhsa --offload-arch=gfx906 \
+// RUN: -fopenmp-offload-mandatory -Xopenmp-target -fexceptions \
+// RUN: -Xopenmp-target -fcxx-exceptions -emit-llvm -S %s -o - 2>&1 | FileCheck %s -v --check-prefix=TEST1
+// TEST1: {{Target 'amdgcn-amd-amdhsa' does not support exception handling. To allow code generation for 'amdgcn-amd-amdhsa', 'try' statements are treated as basic blocks}}
+
+// RUN: %clangxx -O0 -Wall -std=c++17 -fopenmp \
+// RUN: -fopenmp-targets=nvptx64-nvidia-cuda --offload-arch=sm_70 \
+// RUN: -fopenmp-offload-mandatory -Xopenmp-target -fexceptions \
+// RUN: -Xopenmp-target -fcxx-exceptions -emit-llvm -S %s -o - 2>&1 | FileCheck %s -v --check-prefix=TEST2
+// TEST2: {{Target 'nvptx64-nvidia-cuda' does not support exception handling. To allow code generation for 'nvptx64-nvidia-cuda', 'try' statements are treated as basic blocks}}
+
+// Then, we test that exception handling is still allowed on code offloaded to a CPU in TEST3
+// RUN: %clangxx -O0 -Wall -std=c++17 -fopenmp \
+// RUN: -fopenmp-targets=x86_64-pc-linux-gnu --offload-arch=x86_64-pc-linux-gnu \
+// RUN: -fopenmp-offload-mandatory -Xopenmp-target -fexceptions \
+// RUN: -Xopenmp-target -fcxx-exceptions -emit-llvm -S %s -o - 2>&1 | FileCheck %s -v --check-prefix=TEST3
+// TEST3-NOT: {{Target 'x86_64-pc-linux-gnu' does not support exception handling. To allow code generation for 'x86_64-pc-linux-gnu', 'try' statements are treated as basic blocks}}
+
+int main(void) {
+  int retval = -1;
+  #pragma omp target map(always,from:retval)
+  {
+    try {
+      retval = 0;
+    }
+    catch (int e) {
+      retval = e;
+    }
+  }
+  return 0;
+}
\ No newline at end of file
diff --git a/clang/test/OpenMP/target_try_catch_message_fun_call.cpp b/clang/test/OpenMP/target_try_catch_message_fun_call.cpp
new file mode 100644
--- /dev/null
+++ b/clang/test/OpenMP/target_try_catch_message_fun_call.cpp
@@ -0,0 +1,41 @@
+// We first test that we treat 'try' as a basic block on AMD and NVIDIA GPUs in TEST1 and TEST2
+// RUN: %clangxx -O0 -Wall -std=c++17 -fopenmp \
+// RUN: -fopenmp-targets=amdgcn-amd-amdhsa --offload-arch=gfx906 \
+// RUN: -fopenmp-offload-mandatory -Xopenmp-target -fexceptions \
+// RUN: -Xopenmp-target -fcxx-exceptions -emit-llvm -S %s -o - 2>&1 | FileCheck %s -v --check-prefix=TEST1
+// TEST1: {{Target 'amdgcn-amd-amdhsa' does not support exception handling. To allow code generation for 'amdgcn-amd-amdhsa', 'try' statements are treated as basic blocks}}
+
+// RUN: %clangxx -O0 -Wall -std=c++17 -fopenmp \
+// RUN: -fopenmp-targets=nvptx64-nvidia-cuda --offload-arch=sm_70 \
+// RUN: -fopenmp-offload-mandatory -Xopenmp-target -fexceptions \
+// RUN: -Xopenmp-target -fcxx-exceptions -emit-llvm -S %s -o - 2>&1 | FileCheck %s -v --check-prefix=TEST2
+// TEST2: {{Target 'nvptx64-nvidia-cuda' does not support exception handling. To allow code generation for 'nvptx64-nvidia-cuda', 'try' statements are treated as basic blocks}}
+
+// Then, we test that exception handling is still allowed on code offloaded to a CPU in TEST3
+// RUN: %clangxx -O0 -Wall -std=c++17 -fopenmp \
+// RUN: -fopenmp-targets=x86_64-pc-linux-gnu --offload-arch=x86_64-pc-linux-gnu \
+// RUN: -fopenmp-offload-mandatory -Xopenmp-target -fexceptions \
+// RUN: -Xopenmp-target -fcxx-exceptions -emit-llvm -S %s -o - 2>&1 | FileCheck %s -v --check-prefix=TEST3
+// TEST3-NOT: {{Target 'x86_64-pc-linux-gnu' does not support exception handling. To allow code generation for 'x86_64-pc-linux-gnu', 'try' statements are treated as basic blocks}}
+
+#pragma omp declare target
+int foo(){
+ int retval = -1;
+  try {
+    retval = 0;
+  }
+  catch (int e) {
+    retval = e;
+  }
+  return retval;
+}
+#pragma omp end declare target
+
+int main(void) {
+  int retval = -1;
+  #pragma omp target map(always,from:retval)
+  {
+    retval = foo();
+  }
+  return 0;
+}
\ No newline at end of file
diff --git a/clang/test/OpenMP/target_try_catch_throw_message.cpp b/clang/test/OpenMP/target_try_catch_throw_message.cpp
new file mode 100644
--- /dev/null
+++ b/clang/test/OpenMP/target_try_catch_throw_message.cpp
@@ -0,0 +1,36 @@
+// We first test that we treat 'try' as a basic block on AMD and NVIDIA GPUs in TEST1 and TEST2
+// RUN: %clangxx -O0 -Wall -std=c++17 -fopenmp \
+// RUN: -fopenmp-targets=amdgcn-amd-amdhsa --offload-arch=gfx906 \
+// RUN: -fopenmp-offload-mandatory -Xopenmp-target -fexceptions \
+// RUN: -Xopenmp-target -fcxx-exceptions -emit-llvm -S %s -o - 2>&1 | FileCheck %s -v --check-prefix=TEST1
+// TEST1: {{Target 'amdgcn-amd-amdhsa' does not support exception handling. To allow code generation for 'amdgcn-amd-amdhsa', 'try' statements are treated as basic blocks}}
+// TEST1: {{Target 'amdgcn-amd-amdhsa' does not support exception handling. To allow code generation for 'amdgcn-amd-amdhsa', 'throw' expressions will be replaced by traps}}
+
+// RUN: %clangxx -O0 -Wall -std=c++17 -fopenmp \
+// RUN: -fopenmp-targets=nvptx64-nvidia-cuda --offload-arch=sm_70 \
+// RUN: -fopenmp-offload-mandatory -Xopenmp-target -fexceptions \
+// RUN: -Xopenmp-target -fcxx-exceptions -emit-llvm -S %s -o - 2>&1 | FileCheck %s -v --check-prefix=TEST2
+// TEST2: {{Target 'nvptx64-nvidia-cuda' does not support exception handling. To allow code generation for 'nvptx64-nvidia-cuda', 'try' statements are treated as basic blocks}}
+// TEST2: {{Target 'nvptx64-nvidia-cuda' does not support exception handling. To allow code generation for 'nvptx64-nvidia-cuda', 'throw' expressions will be replaced by traps}}
+
+// Then, we test that exception handling is still allowed on code offloaded to a CPU in TEST3
+// RUN: %clangxx -O0 -Wall -std=c++17 -fopenmp \
+// RUN: -fopenmp-targets=x86_64-pc-linux-gnu --offload-arch=x86_64-pc-linux-gnu \
+// RUN: -fopenmp-offload-mandatory -Xopenmp-target -fexceptions \
+// RUN: -Xopenmp-target -fcxx-exceptions -emit-llvm -S %s -o - 2>&1 | FileCheck %s -v --check-prefix=TEST3
+// TEST3-NOT: {{Target 'x86_64-pc-linux-gnu' does not support exception handling. To allow code generation for 'x86_64-pc-linux-gnu', 'try' statements are treated as basic blocks}}
+// TEST3-NOT: {{Target 'x86_64-pc-linux-gnu' does not support exception handling. To allow code generation for 'x86_64-pc-linux-gnu', 'throw' expressions will be replaced by traps}}
+
+int main(void) {
+  int retval = -1;
+  #pragma omp target map(always,from:retval)
+  {
+    try { 
+      throw 404; 
+    }
+    catch (int e) {
+      retval = e;
+    }
+  }
+  return 0;
+}
\ No newline at end of file
diff --git a/clang/test/OpenMP/target_try_catch_throw_message_no_exceptions.cpp b/clang/test/OpenMP/target_try_catch_throw_message_no_exceptions.cpp
new file mode 100644
--- /dev/null
+++ b/clang/test/OpenMP/target_try_catch_throw_message_no_exceptions.cpp
@@ -0,0 +1,22 @@
+// This test makes sure that exception handling is turned off
+// on the device and that the two expected errors appear.
+
+// RUN: %clang_cc1 -x c++ -verify -fopenmp \
+// RUN: -fopenmp-targets=nvptx64-nvidia-cuda %s
+
+// RUN: %clang_cc1 -x c++ -verify -fopenmp \
+// RUN: -fopenmp-targets=amdgcn-amd-amdhsa %s
+
+int main(void) {
+  int retval = -1;
+  #pragma omp target map(always,from:retval)
+  {
+    try { //expected-error{{with exceptions disabled}}
+      throw 404; //expected-error{{with exceptions disabled}}
+    }
+    catch (int e) {
+      retval = e;
+    }
+  }
+  return 0;
+}
\ No newline at end of file