[clang] [llvm] [HIPStdPar] Report reachable C++ exceptions from GPU kernel (PR #228082)
Kunal Dubey via llvm-commits
llvm-commits at lists.llvm.org
Thu Oct 1 06:50:30 PDT 2026
https://github.com/xakep8 created https://github.com/llvm/llvm-project/pull/228082
HipStdPar removes all host functions which are unreachable but there can be cases where some reachable function contains C++ exceptions it is forced to device code and since GPU devices don't support C++ exceptions it must be reported.
This change allows for any reachable C++ exception to be reported as error.
Part of https://github.com/llvm/llvm-project/issues/221941
>From 73f2d82f2248f3cc1e2da67545989844d44e53d1 Mon Sep 17 00:00:00 2001
From: Kunal Dubey <xakep8 at protonmail.com>
Date: Thu, 1 Oct 2026 18:50:35 +0530
Subject: [PATCH] [HIPStdPar] Report reachable C++ exceptions from GPU kernel
HipStdPar removes all host functions which are unreachable but there can
be cases where some reachable function contains C++ exceptions it is
forced to device code and since GPU devices don't support C++ exceptions
it must be reported.
This change allows for any reachable C++ exception to be reported as
error.
---
clang/docs/HIPSupport.md | 7 +-
clang/lib/CodeGen/CGException.cpp | 15 ++
.../unsupported-exceptions.cpp | 139 +++++++++++
llvm/lib/Transforms/HipStdPar/HipStdPar.cpp | 61 ++++-
.../HipStdPar/unsupported-exceptions.ll | 236 ++++++++++++++++++
5 files changed, 446 insertions(+), 12 deletions(-)
create mode 100644 clang/test/CodeGenHipStdPar/unsupported-exceptions.cpp
create mode 100644 llvm/test/Transforms/HipStdPar/unsupported-exceptions.ll
diff --git a/clang/docs/HIPSupport.md b/clang/docs/HIPSupport.md
index e3cfdd4f434274c..e2c9443145117f5 100644
--- a/clang/docs/HIPSupport.md
+++ b/clang/docs/HIPSupport.md
@@ -755,7 +755,10 @@ C++ code:
safely removed in the middle-end.
`CodeGen` is similarly relaxed, with implicitly `__host__` functions being
-emitted as well.
+emitted as well. Because CUDA device code generation otherwise discards the
+exception-handling representation of a C++ `try` statement, HIPStdPar emits an
+unsupported-operation marker that is diagnosed later only if the containing
+function is reachable from an accelerator kernel.
## Implementation - Middle-End
@@ -766,6 +769,8 @@ We add two `opt` passes:
- For all kernels in a `Module`, compute reachability, where a function
`F` is reachable from a kernel `K` if and only if there exists a direct
call-chain rooted in `F` that includes `K`;
+ - Diagnose unsupported constructs, including C++ exception handling, when
+ they are reachable from a kernel;
- Remove all functions that are not reachable from kernels;
- This pass is only run when compiling for the accelerator.
diff --git a/clang/lib/CodeGen/CGException.cpp b/clang/lib/CodeGen/CGException.cpp
index f06edde9739b4a0..26c10727e42d7f8 100644
--- a/clang/lib/CodeGen/CGException.cpp
+++ b/clang/lib/CodeGen/CGException.cpp
@@ -638,6 +638,21 @@ void CodeGenFunction::EmitCXXTryStmt(const CXXTryStmt &S) {
}
void CodeGenFunction::EnterCXXTryStmt(const CXXTryStmt &S, bool IsFnTryBlock) {
+ // HIPStdPar device compilation emits unannotated host functions and removes
+ // the ones that are not reachable from an accelerator kernel in the middle
+ // end. CUDA device code generation otherwise drops the EH representation of
+ // a try statement, so preserving an unsupported-operation marker for the
+ // accelerator code selection pass to diagnose if this function is reachable.
+ if (CGM.getLangOpts().HIPStdPar && CGM.getLangOpts().CUDAIsDevice) {
+ constexpr llvm::StringLiteral MarkerName =
+ "__CXX_EXCEPTION__hipstdpar_unsupported";
+ llvm::FunctionType *MarkerTy =
+ llvm::FunctionType::get(VoidTy, /*isVarArg=*/false);
+ llvm::FunctionCallee Marker =
+ CGM.getModule().getOrInsertFunction(MarkerName, MarkerTy);
+ Builder.CreateCall(Marker);
+ }
+
unsigned NumHandlers = S.getNumHandlers();
EHCatchScope *CatchScope = EHStack.pushCatch(NumHandlers);
diff --git a/clang/test/CodeGenHipStdPar/unsupported-exceptions.cpp b/clang/test/CodeGenHipStdPar/unsupported-exceptions.cpp
new file mode 100644
index 000000000000000..3e7b30170c50d8c
--- /dev/null
+++ b/clang/test/CodeGenHipStdPar/unsupported-exceptions.cpp
@@ -0,0 +1,139 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: not %clang_cc1 -triple amdgcn-amd-amdhsa \
+// RUN: -aux-triple x86_64-unknown-linux-gnu --hipstdpar -x hip \
+// RUN: -fcuda-is-device -fcxx-exceptions -fexceptions \
+// RUN: -mllvm -amdgpu-enable-hipstdpar -emit-llvm -o - -DREACHABLE %s 2>&1 \
+// RUN: | FileCheck %s --check-prefix=REACHABLE
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
+// RUN: -aux-triple x86_64-unknown-linux-gnu --hipstdpar -x hip \
+// RUN: -fcuda-is-device -fcxx-exceptions -fexceptions \
+// RUN: -emit-llvm -o - -DTRY_CATCH %s \
+// RUN: | FileCheck %s --check-prefix=TRY-IR
+// RUN: not %clang_cc1 -triple amdgcn-amd-amdhsa \
+// RUN: -aux-triple x86_64-unknown-linux-gnu --hipstdpar -x hip \
+// RUN: -fcuda-is-device -fcxx-exceptions -fexceptions \
+// RUN: -mllvm -amdgpu-enable-hipstdpar -emit-llvm -o - -DTRY_CATCH %s 2>&1 \
+// RUN: | FileCheck %s --check-prefix=TRY-CATCH
+// RUN: not %clang_cc1 -triple amdgcn-amd-amdhsa \
+// RUN: -aux-triple x86_64-unknown-linux-gnu --hipstdpar -x hip \
+// RUN: -fcuda-is-device -fcxx-exceptions -fexceptions \
+// RUN: -mllvm -amdgpu-enable-hipstdpar -emit-llvm -o - -DFUNCTION_TRY %s 2>&1 \
+// RUN: | FileCheck %s --check-prefix=FUNCTION-TRY
+// RUN: not %clang_cc1 -triple amdgcn-amd-amdhsa \
+// RUN: -aux-triple x86_64-unknown-linux-gnu --hipstdpar -x hip \
+// RUN: -fcuda-is-device -fcxx-exceptions -fexceptions \
+// RUN: -mllvm -amdgpu-enable-hipstdpar -emit-llvm -o - \
+// RUN: -DCONSTRUCTOR_FUNCTION_TRY %s 2>&1 \
+// RUN: | FileCheck %s --check-prefix=CONSTRUCTOR-FUNCTION-TRY
+// RUN: not %clang_cc1 -triple amdgcn-amd-amdhsa \
+// RUN: -aux-triple x86_64-unknown-linux-gnu --hipstdpar -x hip \
+// RUN: -fcuda-is-device -fcxx-exceptions -fexceptions \
+// RUN: -mllvm -amdgpu-enable-hipstdpar -emit-llvm -o - \
+// RUN: -DDESTRUCTOR_FUNCTION_TRY %s 2>&1 \
+// RUN: | FileCheck %s --check-prefix=DESTRUCTOR-FUNCTION-TRY
+// RUN: not %clang_cc1 -triple amdgcn-amd-amdhsa \
+// RUN: -aux-triple x86_64-unknown-linux-gnu --hipstdpar -x hip \
+// RUN: -fcuda-is-device -fcxx-exceptions -fexceptions \
+// RUN: -mllvm -amdgpu-enable-hipstdpar -emit-llvm -o - -DBAD_CAST %s 2>&1 \
+// RUN: | FileCheck %s --check-prefix=BAD-CAST
+// RUN: not %clang_cc1 -triple amdgcn-amd-amdhsa \
+// RUN: -aux-triple x86_64-unknown-linux-gnu --hipstdpar -x hip \
+// RUN: -fcuda-is-device -fcxx-exceptions -fexceptions \
+// RUN: -mllvm -amdgpu-enable-hipstdpar -emit-llvm -o - -DBAD_TYPEID %s 2>&1 \
+// RUN: | FileCheck %s --check-prefix=BAD-TYPEID
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
+// RUN: -aux-triple x86_64-unknown-linux-gnu --hipstdpar -x hip \
+// RUN: -fcuda-is-device -fcxx-exceptions -fexceptions \
+// RUN: -mllvm -amdgpu-enable-hipstdpar -emit-llvm -o - %s \
+// RUN: | FileCheck %s --check-prefix=UNREACHABLE
+
+#define __global__ __attribute__((global))
+
+#if defined(BAD_CAST) || defined(BAD_TYPEID)
+namespace std {
+class type_info;
+}
+
+struct Base {
+ virtual ~Base();
+};
+struct Derived : Base {};
+#endif
+
+#if defined(BAD_CAST)
+Derived &badCast(Base &B) { return dynamic_cast<Derived &>(B); }
+
+__global__ void kernel(Base *B) { (void)badCast(*B); }
+
+// BAD-CAST: error: Accelerator does not support C++ exception handling.
+#elif defined(BAD_TYPEID)
+const std::type_info &badTypeid(Base *B) { return typeid(*B); }
+
+__global__ void kernel(Base *B) { (void)badTypeid(B); }
+
+// BAD-TYPEID: error: Accelerator does not support C++ exception handling.
+#else
+void may_throw();
+
+void throwing_helper() { throw 1; }
+
+void try_helper() {
+ try {
+ may_throw();
+ } catch (...) {
+ }
+}
+
+void function_try_helper() try {
+ may_throw();
+} catch (...) {
+}
+
+struct ConstructorFunctionTry {
+ ConstructorFunctionTry();
+};
+
+ConstructorFunctionTry::ConstructorFunctionTry() try {
+ may_throw();
+} catch (...) {
+}
+
+struct DestructorFunctionTry {
+ ~DestructorFunctionTry();
+};
+
+DestructorFunctionTry::~DestructorFunctionTry() try {
+ may_throw();
+} catch (...) {
+}
+
+__global__ void kernel() {
+#ifdef REACHABLE
+ throwing_helper();
+#elif defined(TRY_CATCH)
+ try_helper();
+#elif defined(FUNCTION_TRY)
+ function_try_helper();
+#elif defined(CONSTRUCTOR_FUNCTION_TRY)
+ ConstructorFunctionTry value;
+#elif defined(DESTRUCTOR_FUNCTION_TRY)
+ DestructorFunctionTry value;
+#endif
+}
+
+// REACHABLE: error: Accelerator does not support C++ exception handling.
+// TRY-CATCH: error: Accelerator does not support C++ exception handling.
+// FUNCTION-TRY: error: Accelerator does not support C++ exception handling.
+// CONSTRUCTOR-FUNCTION-TRY: error: Accelerator does not support C++ exception handling.
+// DESTRUCTOR-FUNCTION-TRY: error: Accelerator does not support C++ exception handling.
+
+// TRY-IR-LABEL: define{{.*}} void @_Z10try_helperv()
+// TRY-IR: call void @__CXX_EXCEPTION__hipstdpar_unsupported()
+
+// UNREACHABLE-NOT: @__cxa_throw
+// UNREACHABLE-NOT: @__CXX_EXCEPTION__hipstdpar_unsupported
+// UNREACHABLE-NOT: @_Z15throwing_helperv
+// UNREACHABLE-NOT: @_Z10try_helperv
+// UNREACHABLE-NOT: @_Z19function_try_helperv
+// UNREACHABLE: define{{.*}} amdgpu_kernel void @_Z6kernelv()
+#endif
diff --git a/llvm/lib/Transforms/HipStdPar/HipStdPar.cpp b/llvm/lib/Transforms/HipStdPar/HipStdPar.cpp
index 17ca146f6a02449..0f4ec6bd30f5e12 100644
--- a/llvm/lib/Transforms/HipStdPar/HipStdPar.cpp
+++ b/llvm/lib/Transforms/HipStdPar/HipStdPar.cpp
@@ -59,6 +59,7 @@
#include "llvm/IR/Constants.h"
#include "llvm/IR/Function.h"
#include "llvm/IR/IRBuilder.h"
+#include "llvm/IR/Instructions.h"
#include "llvm/IR/Intrinsics.h"
#include "llvm/IR/Module.h"
#include "llvm/Transforms/Utils/ModuleUtils.h"
@@ -369,27 +370,62 @@ static inline bool isAcceleratorExecutionRoot(const Function *F) {
return F->getCallingConv() == CallingConv::AMDGPU_KERNEL;
}
-static inline bool checkIfSupported(const Function *F, const CallBase *CB) {
- const auto Dx = F->getName().rfind("__hipstdpar_unsupported");
+static inline bool isCXXExceptionRuntimeFunction(StringRef Name) {
+ return Name == "__cxa_throw" || Name == "__cxa_rethrow" ||
+ Name == "__cxa_bad_cast" || Name == "__cxa_bad_typeid" ||
+ Name == "__cxa_throw_bad_array_new_length" ||
+ Name == "__cxa_rethrow_primary_exception" ||
+ Name == "__cxa_call_unexpected";
+}
- if (Dx == StringRef::npos)
- return true;
+static inline bool checkIfExceptionHandlingIsSupported(const Function *F) {
+ for (const BasicBlock &BB : *F) {
+ for (const Instruction &I : BB) {
+ if (!I.isEHPad() &&
+ !isa<InvokeInst, ResumeInst, CatchReturnInst, CleanupReturnInst>(I))
+ continue;
- const auto N = F->getName().substr(0, Dx);
+ F->getContext().diagnose(DiagnosticInfoUnsupported(
+ *F, "Accelerator does not support C++ exception handling.",
+ I.getDebugLoc(), DS_Error));
+ return false;
+ }
+ }
+
+ return true;
+}
+
+static inline bool checkIfSupported(const Function *F, const CallBase *CB) {
+ StringRef Name = F->getName();
+ const auto Dx = Name.rfind("__hipstdpar_unsupported");
+ // HIPStdPar emits unannotated host functions during device compilation and
+ // removes them here when no kernel can reach them. Defer the unsupported
+ // exception diagnostic until this point for the same reason.
+ const bool IsCXXException = isCXXExceptionRuntimeFunction(Name);
+
+ if (Dx == StringRef::npos && !IsCXXException)
+ return true;
std::string W;
raw_string_ostream OS(W);
- if (N == "__ASM")
- OS << "Accelerator does not support the ASM block:\n"
- << cast<ConstantDataArray>(CB->getArgOperand(0))->getAsCString();
- else
- OS << "Accelerator does not support the " << N << " function.";
+ if (IsCXXException) {
+ OS << "Accelerator does not support C++ exception handling.";
+ } else {
+ const auto N = Name.substr(0, Dx);
+ if (N == "__CXX_EXCEPTION")
+ OS << "Accelerator does not support C++ exception handling.";
+ else if (N == "__ASM")
+ OS << "Accelerator does not support the ASM block:\n"
+ << cast<ConstantDataArray>(CB->getArgOperand(0))->getAsCString();
+ else
+ OS << "Accelerator does not support the " << N << " function.";
+ }
auto Caller = CB->getParent()->getParent();
Caller->getContext().diagnose(
- DiagnosticInfoUnsupported(*Caller, W, CB->getDebugLoc(), DS_Error));
+ DiagnosticInfoUnsupported(*Caller, W, CB->getDebugLoc(), DS_Error));
return false;
}
@@ -411,6 +447,9 @@ PreservedAnalyses
auto F = std::move(Tmp.back());
Tmp.pop_back();
+ if (!checkIfExceptionHandlingIsSupported(F))
+ return PreservedAnalyses::none();
+
for (auto &&N : *CGA[F]) {
if (!N.second)
continue;
diff --git a/llvm/test/Transforms/HipStdPar/unsupported-exceptions.ll b/llvm/test/Transforms/HipStdPar/unsupported-exceptions.ll
new file mode 100644
index 000000000000000..fc90c3c0f39627b
--- /dev/null
+++ b/llvm/test/Transforms/HipStdPar/unsupported-exceptions.ll
@@ -0,0 +1,236 @@
+; RUN: split-file %s %t
+; RUN: not opt -S -mtriple=amdgpu-amd-amdhsa -passes=hipstdpar-select-accelerator-code \
+; RUN: %t/throw.ll 2>&1 | FileCheck %s --check-prefix=THROW
+; RUN: not opt -S -mtriple=amdgpu-amd-amdhsa -passes=hipstdpar-select-accelerator-code \
+; RUN: %t/rethrow.ll 2>&1 | FileCheck %s --check-prefix=RETHROW
+; RUN: not opt -S -mtriple=amdgpu-amd-amdhsa -passes=hipstdpar-select-accelerator-code \
+; RUN: %t/invoke.ll 2>&1 | FileCheck %s --check-prefix=INVOKE
+; RUN: not opt -S -mtriple=amdgpu-amd-amdhsa -passes=hipstdpar-select-accelerator-code \
+; RUN: %t/bad-cast.ll 2>&1 | FileCheck %s --check-prefix=BAD-CAST
+; RUN: not opt -S -mtriple=amdgpu-amd-amdhsa -passes=hipstdpar-select-accelerator-code \
+; RUN: %t/bad-typeid.ll 2>&1 | FileCheck %s --check-prefix=BAD-TYPEID
+; RUN: not opt -S -mtriple=amdgpu-amd-amdhsa -passes=hipstdpar-select-accelerator-code \
+; RUN: %t/bad-array-new.ll 2>&1 | FileCheck %s --check-prefix=BAD-ARRAY-NEW
+; RUN: not opt -S -mtriple=amdgpu-amd-amdhsa -passes=hipstdpar-select-accelerator-code \
+; RUN: %t/rethrow-primary.ll 2>&1 | FileCheck %s --check-prefix=RETHROW-PRIMARY
+; RUN: not opt -S -mtriple=amdgpu-amd-amdhsa -passes=hipstdpar-select-accelerator-code \
+; RUN: %t/call-unexpected.ll 2>&1 | FileCheck %s --check-prefix=CALL-UNEXPECTED
+; RUN: not opt -S -mtriple=amdgpu-amd-amdhsa -passes=hipstdpar-select-accelerator-code \
+; RUN: %t/catch-only.ll 2>&1 | FileCheck %s --check-prefix=CATCH-ONLY
+; RUN: not opt -S -mtriple=amdgpu-amd-amdhsa -passes=hipstdpar-select-accelerator-code \
+; RUN: %t/try-marker.ll 2>&1 | FileCheck %s --check-prefix=TRY-MARKER
+; RUN: opt -S -mtriple=amdgpu-amd-amdhsa -passes=hipstdpar-select-accelerator-code \
+; RUN: %t/unreachable.ll | FileCheck %s --check-prefix=UNREACHABLE
+; RUN: opt -S -mtriple=amdgpu-amd-amdhsa -passes=hipstdpar-select-accelerator-code \
+; RUN: %t/similar-names.ll | FileCheck %s --check-prefix=SIMILAR-NAMES
+
+; THROW: error: {{.*}} in function throwing_helper void (): Accelerator does not support C++ exception handling.
+; RETHROW: error: {{.*}} in function rethrowing_helper void (): Accelerator does not support C++ exception handling.
+; INVOKE: error: {{.*}} in function kernel void (): Accelerator does not support C++ exception handling.
+; BAD-CAST: error: {{.*}} in function kernel void (): Accelerator does not support C++ exception handling.
+; BAD-TYPEID: error: {{.*}} in function kernel void (): Accelerator does not support C++ exception handling.
+; BAD-ARRAY-NEW: error: {{.*}} in function kernel void (): Accelerator does not support C++ exception handling.
+; RETHROW-PRIMARY: error: {{.*}} in function kernel void (): Accelerator does not support C++ exception handling.
+; CALL-UNEXPECTED: error: {{.*}} in function kernel void (): Accelerator does not support C++ exception handling.
+; CATCH-ONLY: error: {{.*}} in function kernel void (): Accelerator does not support C++ exception handling.
+; TRY-MARKER: error: {{.*}} in function helper void (): Accelerator does not support C++ exception handling.
+
+; UNREACHABLE-NOT: @host_only
+; UNREACHABLE-NOT: @host_catch_only
+; UNREACHABLE-NOT: @__cxa_
+; UNREACHABLE-NOT: @__CXX_EXCEPTION__hipstdpar_unsupported
+; UNREACHABLE: define amdgpu_kernel void @kernel()
+; UNREACHABLE-NOT: @host_only
+; UNREACHABLE-NOT: @host_catch_only
+; UNREACHABLE-NOT: @__cxa_
+; UNREACHABLE-NOT: @__CXX_EXCEPTION__hipstdpar_unsupported
+
+; SIMILAR-NAMES: define amdgpu_kernel void @kernel()
+; SIMILAR-NAMES: call void @__cxa_throwing()
+; SIMILAR-NAMES: call void @my___cxa_rethrow()
+; SIMILAR-NAMES: call ptr @__cxa_allocate_exception(i64 4)
+
+;--- throw.ll
+define void @throwing_helper() {
+entry:
+ call void @__cxa_throw(ptr null, ptr null, ptr null)
+ unreachable
+}
+
+define amdgpu_kernel void @kernel() {
+entry:
+ call void @throwing_helper()
+ ret void
+}
+
+declare void @__cxa_throw(ptr, ptr, ptr)
+
+;--- rethrow.ll
+define void @rethrowing_helper() {
+entry:
+ call void @__cxa_rethrow()
+ unreachable
+}
+
+define amdgpu_kernel void @kernel() {
+entry:
+ call void @rethrowing_helper()
+ ret void
+}
+
+declare void @__cxa_rethrow()
+
+;--- invoke.ll
+define amdgpu_kernel void @kernel() personality ptr @__gxx_personality_v0 {
+entry:
+ invoke void @__cxa_throw(ptr null, ptr null, ptr null)
+ to label %normal unwind label %cleanup
+
+normal:
+ unreachable
+
+cleanup:
+ %landing = landingpad { ptr, i32 }
+ cleanup
+ resume { ptr, i32 } %landing
+}
+
+declare void @__cxa_throw(ptr, ptr, ptr)
+declare i32 @__gxx_personality_v0(...)
+
+;--- bad-cast.ll
+define amdgpu_kernel void @kernel() {
+entry:
+ call void @__cxa_bad_cast()
+ unreachable
+}
+
+declare void @__cxa_bad_cast()
+
+;--- bad-typeid.ll
+define amdgpu_kernel void @kernel() {
+entry:
+ call void @__cxa_bad_typeid()
+ unreachable
+}
+
+declare void @__cxa_bad_typeid()
+
+;--- bad-array-new.ll
+define amdgpu_kernel void @kernel() {
+entry:
+ call void @__cxa_throw_bad_array_new_length()
+ unreachable
+}
+
+declare void @__cxa_throw_bad_array_new_length()
+
+;--- rethrow-primary.ll
+define amdgpu_kernel void @kernel() {
+entry:
+ call void @__cxa_rethrow_primary_exception(ptr null)
+ unreachable
+}
+
+declare void @__cxa_rethrow_primary_exception(ptr)
+
+;--- call-unexpected.ll
+define amdgpu_kernel void @kernel() {
+entry:
+ call void @__cxa_call_unexpected(ptr null)
+ unreachable
+}
+
+declare void @__cxa_call_unexpected(ptr)
+
+;--- catch-only.ll
+define amdgpu_kernel void @kernel() personality ptr @__gxx_personality_v0 {
+entry:
+ invoke void @may_throw()
+ to label %exit unwind label %catch
+
+catch:
+ %landing = landingpad { ptr, i32 }
+ catch ptr null
+ ret void
+
+exit:
+ ret void
+}
+
+declare void @may_throw()
+declare i32 @__gxx_personality_v0(...)
+
+;--- try-marker.ll
+define void @helper() {
+entry:
+ call void @__CXX_EXCEPTION__hipstdpar_unsupported()
+ call void @may_throw()
+ ret void
+}
+
+define amdgpu_kernel void @kernel() {
+entry:
+ call void @helper()
+ ret void
+}
+
+declare void @__CXX_EXCEPTION__hipstdpar_unsupported()
+declare void @may_throw()
+
+;--- unreachable.ll
+define void @host_only() {
+entry:
+ call void @__CXX_EXCEPTION__hipstdpar_unsupported()
+ call void @__cxa_throw(ptr null, ptr null, ptr null)
+ call void @__cxa_rethrow()
+ call void @__cxa_bad_cast()
+ call void @__cxa_bad_typeid()
+ call void @__cxa_throw_bad_array_new_length()
+ call void @__cxa_rethrow_primary_exception(ptr null)
+ call void @__cxa_call_unexpected(ptr null)
+ unreachable
+}
+
+define void @host_catch_only() personality ptr @__gxx_personality_v0 {
+entry:
+ invoke void @may_throw()
+ to label %exit unwind label %catch
+
+catch:
+ %landing = landingpad { ptr, i32 }
+ catch ptr null
+ ret void
+
+exit:
+ ret void
+}
+
+define amdgpu_kernel void @kernel() {
+entry:
+ ret void
+}
+
+declare void @__cxa_throw(ptr, ptr, ptr)
+declare void @__cxa_rethrow()
+declare void @__cxa_bad_cast()
+declare void @__cxa_bad_typeid()
+declare void @__cxa_throw_bad_array_new_length()
+declare void @__cxa_rethrow_primary_exception(ptr)
+declare void @__cxa_call_unexpected(ptr)
+declare void @__CXX_EXCEPTION__hipstdpar_unsupported()
+declare void @may_throw()
+declare i32 @__gxx_personality_v0(...)
+
+;--- similar-names.ll
+define amdgpu_kernel void @kernel() {
+entry:
+ call void @__cxa_throwing()
+ call void @my___cxa_rethrow()
+ %exception = call ptr @__cxa_allocate_exception(i64 4)
+ ret void
+}
+
+declare void @__cxa_throwing()
+declare void @my___cxa_rethrow()
+declare ptr @__cxa_allocate_exception(i64)
More information about the llvm-commits
mailing list