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 <[email protected]> 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) _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
