Author: Anton Rydahl Date: 2023-08-29T15:05:59-07:00 New Revision: 4c62e943b7178127861ca39163a0ed4caeb14943
URL: https://github.com/llvm/llvm-project/commit/4c62e943b7178127861ca39163a0ed4caeb14943 DIFF: https://github.com/llvm/llvm-project/commit/4c62e943b7178127861ca39163a0ed4caeb14943.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 ``` 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 ``` ./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..d0104247de98ae --- /dev/null +++ b/clang/test/OpenMP/amdgpu_exceptions.cpp @@ -0,0 +1,48 @@ +// REQUIRES: amdgpu-registered-target + +/** + * 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 00000000000000..9afa7261a511b2 --- /dev/null +++ b/clang/test/OpenMP/amdgpu_throw.cpp @@ -0,0 +1,40 @@ +// REQUIRES: amdgpu-registered-target + +/** + * 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 00000000000000..f1152f78bec192 --- /dev/null +++ b/clang/test/OpenMP/amdgpu_throw_trap.cpp @@ -0,0 +1,13 @@ +// REQUIRES: amdgpu-registered-target + +// 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..b82656400270d5 --- /dev/null +++ b/clang/test/OpenMP/amdgpu_try_catch.cpp @@ -0,0 +1,47 @@ +// REQUIRES: amdgpu-registered-target + +/** + * 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 00000000000000..89f62dc33e47a4 --- /dev/null +++ b/clang/test/OpenMP/nvptx_exceptions.cpp @@ -0,0 +1,48 @@ +// REQUIRES: nvptx-registered-target + +/** + * 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 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@-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..f17f75b401007b --- /dev/null +++ b/clang/test/OpenMP/nvptx_throw.cpp @@ -0,0 +1,40 @@ +// REQUIRES: nvptx-registered-target + +/** + * 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 00000000000000..c1c76c4e1b18c9 --- /dev/null +++ b/clang/test/OpenMP/nvptx_throw_trap.cpp @@ -0,0 +1,13 @@ +// REQUIRES: nvptx-registered-target + +// 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..c57aa91d851a4c --- /dev/null +++ b/clang/test/OpenMP/nvptx_try_catch.cpp @@ -0,0 +1,47 @@ +// REQUIRES: nvptx-registered-target + +/** + * 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 00000000000000..effa76f0016dbe --- /dev/null +++ b/clang/test/OpenMP/x86_target_exceptions.cpp @@ -0,0 +1,16 @@ +// REQUIRES: x86-registered-target + +// 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..b55775287daa69 --- /dev/null +++ b/clang/test/OpenMP/x86_target_throw.cpp @@ -0,0 +1,9 @@ +// REQUIRES: x86-registered-target + +// 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..f18e6aaf99f90a --- /dev/null +++ b/clang/test/OpenMP/x86_target_try_catch.cpp @@ -0,0 +1,16 @@ +// REQUIRES: x86-registered-target + +// 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 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
