[clang] 0cfc2db - [OpenMP] Allow exceptions in target regions when offloading to GPUs
via cfe-commits
cfe-commits at lists.llvm.org
Mon Aug 28 22:36:44 PDT 2023
Author: Anton Rydahl
Date: 2023-08-28T22:36:13-07:00
New Revision: 0cfc2dba93b172802b580713a492ea14148a0218
URL: https://github.com/llvm/llvm-project/commit/0cfc2dba93b172802b580713a492ea14148a0218
DIFF: https://github.com/llvm/llvm-project/commit/0cfc2dba93b172802b580713a492ea14148a0218.diff
LOG: [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
Added:
clang/test/OpenMP/amdgpu_exceptions.cpp
clang/test/OpenMP/amdgpu_throw.cpp
clang/test/OpenMP/amdgpu_throw_trap.cpp
clang/test/OpenMP/amdgpu_try_catch.cpp
clang/test/OpenMP/nvptx_exceptions.cpp
clang/test/OpenMP/nvptx_throw.cpp
clang/test/OpenMP/nvptx_throw_trap.cpp
clang/test/OpenMP/nvptx_try_catch.cpp
clang/test/OpenMP/x86_target_exceptions.cpp
clang/test/OpenMP/x86_target_throw.cpp
clang/test/OpenMP/x86_target_try_catch.cpp
Modified:
clang/include/clang/Basic/DiagnosticCommonKinds.td
clang/include/clang/Basic/DiagnosticGroups.td
clang/lib/CodeGen/CGException.cpp
clang/lib/Sema/SemaExprCXX.cpp
clang/lib/Sema/SemaStmt.cpp
clang/test/OpenMP/nvptx_target_exceptions_messages.cpp
Removed:
################################################################################
diff --git a/clang/include/clang/Basic/DiagnosticCommonKinds.td b/clang/include/clang/Basic/DiagnosticCommonKinds.td
index cd72e254ea3b1a..f2df283c74829f 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_
diff ers : Error<
"options %0 and %1 are set to
diff erent 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 d1aa51393ef357..00c458fb23e73e 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 9cb7d4c7731deb..3996f2948349cb 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 a6e0a8abf81b9a..472fbdbdb5d0e6 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 70a549938d080d..88ea3f1c3349dd 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 00000000000000..f381ec30b45a02
--- /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 at 37 {{cannot use 'try' with exceptions disabled}}
+// noexceptions-error at 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 00000000000000..bccba270d05bc0
--- /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 at 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 00000000000000..513b543ccaa8f4
--- /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 00000000000000..7aa20419edcd45
--- /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 at 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 00000000000000..e59fb1f3d30bc2
--- /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 at 37 {{cannot use 'try' with exceptions disabled}}
+// noexceptions-error at 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 9f267f5606e2d7..5d1d46cadff260 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 at -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 00000000000000..6d5774551fe94e
--- /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 at 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 00000000000000..d09f4dfd885714
--- /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 00000000000000..c2a0786aefd539
--- /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 at 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 00000000000000..bc5e485b88345b
--- /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 00000000000000..03fef3faa83d94
--- /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 00000000000000..e878d2ee0929b7
--- /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
More information about the cfe-commits
mailing list