[OpenMP] Allow exceptions in target regions when offloading to GPUs

The motivation for this patch is that many code bases use exception handling. As GPUs are not expected to support exception handling in the near future, we can experiment with compiling the code for GPU targets anyway. This will
allow us to run the code, as long as no exception is thrown.

The overall idea is very simple:
- If a throw expression is compiled to AMDGCN or NVPTX, it is replaced with a trap during code generation.
- If a try/catch statement is compiled to AMDGCN or AMDHSA, we ganerate code for the try statement as if it were a basic block.

With this patch, the compilation of the following example
```{C++}
int gaussian_sum(int a,int b){
	if ((a + b) % 2 == 0) {throw -1;};
	return (a+b) * ((a+b)/2);
}

int main(void) {
	int gauss = 0;
	#pragma omp target map(from:gauss)
	{
		try {
			gauss = gaussian_sum(1,100);
		}
		catch (int e){
			gauss = e;
		}
	}
	std::cout << "GaussianSum(1,100)="<<gauss<<std::endl;
        #pragma omp target map(from:gauss)
        {
                try {
                     	gauss = gaussian_sum(1,101);
                }
                catch (int e){
                        gauss = e;
                }
        }
	std::cout << "GaussianSum(1,101)="<<gauss<<std::endl;
	return (gauss > 1) ? 0 : 1;
}
```
with offloading to `gfx906` results in
```{bash}
./bin/target_try_minimal_fail
GaussianSum(1,100)=5050
AMDGPU fatal error 1: Received error in queue 0x155555506000: HSA_STATUS_ERROR_EXCEPTION: An HSAIL operation resulted in a hardware exception.
zsh: abort (core dumped)
```

Reviewed By: jdoerfert

Differential Revision: https://reviews.llvm.org/D153924

Author: AntonRydahl

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>;
 }

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.

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) {

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)

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)

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

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

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

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

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

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;

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