diff options
Diffstat (limited to 'libgomp/testsuite/libgomp.c++')
25 files changed, 750 insertions, 2 deletions
diff --git a/libgomp/testsuite/libgomp.c++/allocator-1.C b/libgomp/testsuite/libgomp.c++/allocator-1.C new file mode 100644 index 0000000..49425386 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/allocator-1.C @@ -0,0 +1,171 @@ +// { dg-do run } + +#include <omp.h> +#include <memory> +#include <limits> + +template<typename T, template<typename> class Alloc> +void test (T const initial_value = T()) +{ + using Allocator = Alloc<T>; + Allocator a; + using Traits = std::allocator_traits<Allocator>; + static_assert (__is_same(typename Traits::allocator_type, Allocator )); + static_assert (__is_same(typename Traits::value_type, T )); + static_assert (__is_same(typename Traits::pointer, T* )); + static_assert (__is_same(typename Traits::const_pointer, T const* )); + static_assert (__is_same(typename Traits::void_pointer, void* )); + static_assert (__is_same(typename Traits::const_void_pointer, void const* )); + static_assert (__is_same(typename Traits::difference_type, __PTRDIFF_TYPE__)); + static_assert (__is_same(typename Traits::size_type, __SIZE_TYPE__ )); + static_assert (Traits::propagate_on_container_copy_assignment::value == false); + static_assert (Traits::propagate_on_container_move_assignment::value == false); + static_assert (Traits::propagate_on_container_swap::value == false); + static_assert (Traits::is_always_equal::value == true); + + static constexpr __SIZE_TYPE__ correct_max_size + = std::numeric_limits<__SIZE_TYPE__>::max () / sizeof (T); + if (Traits::max_size (a) != correct_max_size) + __builtin_abort (); + + static constexpr __SIZE_TYPE__ alloc_count = 1; + T *p = Traits::allocate (a, alloc_count); + if (p == nullptr) + __builtin_abort (); + Traits::construct (a, p, initial_value); + if (*p != initial_value) + __builtin_abort (); + Traits::destroy (a, p); + Traits::deallocate (a, p, alloc_count); + /* Not interesting but might as well test it. */ + static_cast<void>(Traits::select_on_container_copy_construction (a)); + + if (!(a == Allocator())) + __builtin_abort (); + if (a != Allocator()) + __builtin_abort (); + if (!(a == Alloc<void>())) + __builtin_abort (); + if (a != Alloc<void>()) + __builtin_abort (); +} + +#define CHECK_INEQUALITY(other_alloc_templ, type) \ +do { \ + /* Skip tests for itself, those are equal. Intantiate each */ \ + /* one with void so we can easily tell if they are the same. */ \ + if (!__is_same (AllocTempl<void>, other_alloc_templ<void>)) \ + { \ + other_alloc_templ<type> other; \ + if (a == other) \ + __builtin_abort (); \ + if (!(a != other)) \ + __builtin_abort (); \ + } \ +} while (false) + +template<typename T, template<typename> class AllocTempl> +void test_inequality () +{ + using Allocator = AllocTempl<T>; + Allocator a; + CHECK_INEQUALITY (omp::allocator::null_allocator, void); + CHECK_INEQUALITY (omp::allocator::default_mem, void); + CHECK_INEQUALITY (omp::allocator::large_cap_mem, void); + CHECK_INEQUALITY (omp::allocator::const_mem, void); + CHECK_INEQUALITY (omp::allocator::high_bw_mem, void); + CHECK_INEQUALITY (omp::allocator::low_lat_mem, void); + CHECK_INEQUALITY (omp::allocator::cgroup_mem, void); + CHECK_INEQUALITY (omp::allocator::pteam_mem, void); + CHECK_INEQUALITY (omp::allocator::thread_mem, void); +#ifdef __gnu_linux__ + /* Pinning not implemented on other targets. */ + CHECK_INEQUALITY (ompx::allocator::gnu_pinned_mem, void); +#endif + /* And again with the same type passed to the allocator. */ + CHECK_INEQUALITY (omp::allocator::null_allocator, T); + CHECK_INEQUALITY (omp::allocator::default_mem, T); + CHECK_INEQUALITY (omp::allocator::large_cap_mem, T); + CHECK_INEQUALITY (omp::allocator::const_mem, T); + CHECK_INEQUALITY (omp::allocator::high_bw_mem, T); + CHECK_INEQUALITY (omp::allocator::low_lat_mem, T); + CHECK_INEQUALITY (omp::allocator::cgroup_mem, T); + CHECK_INEQUALITY (omp::allocator::pteam_mem, T); + CHECK_INEQUALITY (omp::allocator::thread_mem, T); +#ifdef __gnu_linux__ + CHECK_INEQUALITY (ompx::allocator::gnu_pinned_mem, T); +#endif +} + +#undef CHECK_INEQUALITY + +struct S +{ + int _v0; + bool _v1; + float _v2; + + bool operator== (S const& other) const noexcept { + return _v0 == other._v0 + && _v1 == other._v1 + && _v2 == other._v2; + } + bool operator!= (S const& other) const noexcept { + return !this->operator==(other); + } +}; + +int main () +{ + test<int, omp::allocator::null_allocator>(42); + test<int, omp::allocator::default_mem>(42); + test<int, omp::allocator::large_cap_mem>(42); + test<int, omp::allocator::const_mem>(42); + test<int, omp::allocator::high_bw_mem>(42); + test<int, omp::allocator::low_lat_mem>(42); + test<int, omp::allocator::cgroup_mem>(42); + test<int, omp::allocator::pteam_mem>(42); + test<int, omp::allocator::thread_mem>(42); +#ifdef __gnu_linux__ + test<int, ompx::allocator::gnu_pinned_mem>(42); +#endif + + test<long long, omp::allocator::null_allocator>(42); + test<long long, omp::allocator::default_mem>(42); + test<long long, omp::allocator::large_cap_mem>(42); + test<long long, omp::allocator::const_mem>(42); + test<long long, omp::allocator::high_bw_mem>(42); + test<long long, omp::allocator::low_lat_mem>(42); + test<long long, omp::allocator::cgroup_mem>(42); + test<long long, omp::allocator::pteam_mem>(42); + test<long long, omp::allocator::thread_mem>(42); +#ifdef __gnu_linux__ + test<long long, ompx::allocator::gnu_pinned_mem>(42); +#endif + + test<S, omp::allocator::null_allocator>( S{42, true, 128.f}); + test<S, omp::allocator::default_mem>( S{42, true, 128.f}); + test<S, omp::allocator::large_cap_mem>( S{42, true, 128.f}); + test<S, omp::allocator::const_mem>( S{42, true, 128.f}); + test<S, omp::allocator::high_bw_mem>( S{42, true, 128.f}); + test<S, omp::allocator::low_lat_mem>( S{42, true, 128.f}); + test<S, omp::allocator::cgroup_mem>( S{42, true, 128.f}); + test<S, omp::allocator::pteam_mem>( S{42, true, 128.f}); + test<S, omp::allocator::thread_mem>( S{42, true, 128.f}); +#ifdef __gnu_linux__ + test<S, ompx::allocator::gnu_pinned_mem>(S{42, true, 128.f}); +#endif + + test_inequality<int, omp::allocator::null_allocator>(); + test_inequality<int, omp::allocator::default_mem>(); + test_inequality<int, omp::allocator::large_cap_mem>(); + test_inequality<int, omp::allocator::const_mem>(); + test_inequality<int, omp::allocator::high_bw_mem>(); + test_inequality<int, omp::allocator::low_lat_mem>(); + test_inequality<int, omp::allocator::cgroup_mem>(); + test_inequality<int, omp::allocator::pteam_mem>(); + test_inequality<int, omp::allocator::thread_mem>(); +#ifdef __gnu_linux__ + test_inequality<int, ompx::allocator::gnu_pinned_mem>(); +#endif +} diff --git a/libgomp/testsuite/libgomp.c++/allocator-2.C b/libgomp/testsuite/libgomp.c++/allocator-2.C new file mode 100644 index 0000000..ca94fc7 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/allocator-2.C @@ -0,0 +1,141 @@ +// { dg-do run } +// { dg-additional-options "-Wno-psabi" } + +#include <omp.h> +#include <vector> + +template<typename T> +bool ptr_is_aligned(T *ptr, std::size_t alignment) +{ + /* ALIGNMENT must be a power of 2. */ + if ((alignment & (alignment - 1)) != 0) + __builtin_abort (); + __UINTPTR_TYPE__ ptr_value + = reinterpret_cast<__UINTPTR_TYPE__>(static_cast<void*>(ptr)); + return (ptr_value % alignment) == 0; +} + +template<typename T, template<typename> class Alloc> +void f (T v0, T v1, T v2, T v3) +{ + std::vector<T, Alloc<T>> vec; + vec.push_back (v0); + vec.push_back (v1); + vec.push_back (v2); + vec.push_back (v3); + if (vec.at (0) != v0) + __builtin_abort (); + if (vec.at (1) != v1) + __builtin_abort (); + if (vec.at (2) != v2) + __builtin_abort (); + if (vec.at (3) != v3) + __builtin_abort (); + if (!ptr_is_aligned (&vec.at (0), alignof (T))) + __builtin_abort (); + if (!ptr_is_aligned (&vec.at (1), alignof (T))) + __builtin_abort (); + if (!ptr_is_aligned (&vec.at (2), alignof (T))) + __builtin_abort (); + if (!ptr_is_aligned (&vec.at (3), alignof (T))) + __builtin_abort (); +} + +struct S0 +{ + int _v0; + bool _v1; + float _v2; + + bool operator== (S0 const& other) const noexcept { + return _v0 == other._v0 + && _v1 == other._v1 + && _v2 == other._v2; + } + bool operator!= (S0 const& other) const noexcept { + return !this->operator==(other); + } +}; + +struct alignas(128) S1 +{ + int _v0; + bool _v1; + float _v2; + + bool operator== (S1 const& other) const noexcept { + return _v0 == other._v0 + && _v1 == other._v1 + && _v2 == other._v2; + } + bool operator!= (S1 const& other) const noexcept { + return !this->operator==(other); + } +}; + +/* Note: the test for const_mem should be disabled in the future. */ + +int main () +{ + f<int, omp::allocator::null_allocator >(0, 1, 2, 3); + f<int, omp::allocator::default_mem >(0, 1, 2, 3); + f<int, omp::allocator::large_cap_mem >(0, 1, 2, 3); + f<int, omp::allocator::const_mem >(0, 1, 2, 3); + f<int, omp::allocator::high_bw_mem >(0, 1, 2, 3); + f<int, omp::allocator::low_lat_mem >(0, 1, 2, 3); + f<int, omp::allocator::cgroup_mem >(0, 1, 2, 3); + f<int, omp::allocator::pteam_mem >(0, 1, 2, 3); + f<int, omp::allocator::thread_mem >(0, 1, 2, 3); +#ifdef __gnu_linux__ + /* Pinning not implemented on other targets. */ + f<int, ompx::allocator::gnu_pinned_mem>(0, 1, 2, 3); +#endif + + f<long long, omp::allocator::null_allocator >(0, 1, 2, 3); + f<long long, omp::allocator::default_mem >(0, 1, 2, 3); + f<long long, omp::allocator::large_cap_mem >(0, 1, 2, 3); + f<long long, omp::allocator::const_mem >(0, 1, 2, 3); + f<long long, omp::allocator::high_bw_mem >(0, 1, 2, 3); + f<long long, omp::allocator::low_lat_mem >(0, 1, 2, 3); + f<long long, omp::allocator::cgroup_mem >(0, 1, 2, 3); + f<long long, omp::allocator::pteam_mem >(0, 1, 2, 3); + f<long long, omp::allocator::thread_mem >(0, 1, 2, 3); +#ifdef __gnu_linux__ + f<long long, ompx::allocator::gnu_pinned_mem>(0, 1, 2, 3); +#endif + + S0 s0_0{ 42, true, 111128.f}; + S0 s0_1{ 142, false, 11128.f}; + S0 s0_2{ 1142, true, 1128.f}; + S0 s0_3{11142, false, 128.f}; + f<S0, omp::allocator::null_allocator >(s0_0, s0_1, s0_2, s0_3); + f<S0, omp::allocator::default_mem >(s0_0, s0_1, s0_2, s0_3); + f<S0, omp::allocator::large_cap_mem >(s0_0, s0_1, s0_2, s0_3); + f<S0, omp::allocator::const_mem >(s0_0, s0_1, s0_2, s0_3); + f<S0, omp::allocator::high_bw_mem >(s0_0, s0_1, s0_2, s0_3); + f<S0, omp::allocator::low_lat_mem >(s0_0, s0_1, s0_2, s0_3); + f<S0, omp::allocator::cgroup_mem >(s0_0, s0_1, s0_2, s0_3); + f<S0, omp::allocator::pteam_mem >(s0_0, s0_1, s0_2, s0_3); + f<S0, omp::allocator::thread_mem >(s0_0, s0_1, s0_2, s0_3); +#ifdef __gnu_linux__ + f<S0, ompx::allocator::gnu_pinned_mem>(s0_0, s0_1, s0_2, s0_3); +#endif + + S1 s1_0{ 42, true, 111128.f}; + S1 s1_1{ 142, false, 11128.f}; + S1 s1_2{ 1142, true, 1128.f}; + S1 s1_3{11142, false, 128.f}; + + f<S1, omp::allocator::null_allocator >(s1_0, s1_1, s1_2, s1_3); + f<S1, omp::allocator::default_mem >(s1_0, s1_1, s1_2, s1_3); + f<S1, omp::allocator::large_cap_mem >(s1_0, s1_1, s1_2, s1_3); + f<S1, omp::allocator::const_mem >(s1_0, s1_1, s1_2, s1_3); + f<S1, omp::allocator::high_bw_mem >(s1_0, s1_1, s1_2, s1_3); + f<S1, omp::allocator::low_lat_mem >(s1_0, s1_1, s1_2, s1_3); + f<S1, omp::allocator::cgroup_mem >(s1_0, s1_1, s1_2, s1_3); + f<S1, omp::allocator::pteam_mem >(s1_0, s1_1, s1_2, s1_3); + f<S1, omp::allocator::thread_mem >(s1_0, s1_1, s1_2, s1_3); +#ifdef __gnu_linux__ + f<S1, ompx::allocator::gnu_pinned_mem>(s1_0, s1_1, s1_2, s1_3); +#endif +} diff --git a/libgomp/testsuite/libgomp.c++/pr106445-1-O0.C b/libgomp/testsuite/libgomp.c++/pr106445-1-O0.C new file mode 100644 index 0000000..bcd499c --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/pr106445-1-O0.C @@ -0,0 +1,3 @@ +// { dg-additional-options -O0 } + +#include "pr106445-1.C" diff --git a/libgomp/testsuite/libgomp.c++/pr106445-1.C b/libgomp/testsuite/libgomp.c++/pr106445-1.C new file mode 100644 index 0000000..329ce62 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/pr106445-1.C @@ -0,0 +1,18 @@ +#include <vector> + +int main() +{ +#pragma omp target + { + { + std::vector<int> v; + if (!v.empty()) + __builtin_abort(); + } + { + std::vector<int> v(100); + if (v.capacity() < 100) + __builtin_abort(); + } + } +} diff --git a/libgomp/testsuite/libgomp.c++/pr119692-1-1.C b/libgomp/testsuite/libgomp.c++/pr119692-1-1.C new file mode 100644 index 0000000..1f59b15 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/pr119692-1-1.C @@ -0,0 +1,10 @@ +/* PR119692 "C++ 'typeinfo', 'vtable' vs. OpenACC, OpenMP 'target' offloading" */ + +/* { dg-additional-options -UDEFAULT } + Wrong code for offloading execution. + { dg-additional-options -foffload=disable } */ +/* { dg-additional-options -fdump-tree-gimple } */ + +#include "../libgomp.oacc-c++/pr119692-1-1.C" + +/* { dg-final { scan-tree-dump-not {(?n)#pragma omp target .* map\(tofrom:_ZTI2C2 \[len: [0-9]+\] \[runtime_implicit\]\) map\(tofrom:_ZTI2C1 \[len: [0-9]+\] \[runtime_implicit\]\) map\(tofrom:_ZTV2C1 \[len: [0-9]+\] \[runtime_implicit\]\)$} gimple { xfail *-*-* } } } */ diff --git a/libgomp/testsuite/libgomp.c++/pr119692-1-2.C b/libgomp/testsuite/libgomp.c++/pr119692-1-2.C new file mode 100644 index 0000000..e7ac818 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/pr119692-1-2.C @@ -0,0 +1,11 @@ +/* PR119692 "C++ 'typeinfo', 'vtable' vs. OpenACC, OpenMP 'target' offloading" */ + +/* { dg-additional-options -DDEFAULT=defaultmap(none) } + Fails to compile. + { dg-do compile } */ + +#include "pr119692-1-1.C" + +/* { dg-bogus {error: '_ZTV2C1' not specified in enclosing 'target'} PR119692 { xfail *-*-* } 0 } + { dg-bogus {error: '_ZTI2C2' not specified in enclosing 'target'} PR119692 { xfail *-*-* } 0 } + { dg-bogus {error: '_ZTI2C1' not specified in enclosing 'target'} PR119692 { xfail *-*-* } 0 } */ diff --git a/libgomp/testsuite/libgomp.c++/pr119692-1-3.C b/libgomp/testsuite/libgomp.c++/pr119692-1-3.C new file mode 100644 index 0000000..733feb8 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/pr119692-1-3.C @@ -0,0 +1,10 @@ +/* PR119692 "C++ 'typeinfo', 'vtable' vs. OpenACC, OpenMP 'target' offloading" */ + +/* { dg-additional-options -DDEFAULT=defaultmap(present) } + Wrong code for offloading execution. + { dg-xfail-run-if PR119692 { offload_device } } */ +/* { dg-additional-options -fdump-tree-gimple } */ + +#include "pr119692-1-1.C" + +/* { dg-final { scan-tree-dump-not {(?n)#pragma omp target .* defaultmap\(present\) map\(force_present:_ZTI2C2 \[len: [0-9]+\] \[runtime_implicit\]\) map\(force_present:_ZTI2C1 \[len: [0-9]+\] \[runtime_implicit\]\) map\(force_present:_ZTV2C1 \[len: [0-9]+\] \[runtime_implicit\]\)$} gimple { xfail *-*-* } } } */ diff --git a/libgomp/testsuite/libgomp.c++/pr119692-1-4.C b/libgomp/testsuite/libgomp.c++/pr119692-1-4.C new file mode 100644 index 0000000..6995f26 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/pr119692-1-4.C @@ -0,0 +1,10 @@ +/* PR119692 "C++ 'typeinfo', 'vtable' vs. OpenACC, OpenMP 'target' offloading" */ + +/* { dg-additional-options -DDEFAULT=defaultmap(firstprivate) } + Wrong code for offloading execution. + { dg-xfail-run-if PR119692 { offload_device } } */ +/* { dg-additional-options -fdump-tree-gimple } */ + +#include "pr119692-1-1.C" + +/* { dg-final { scan-tree-dump-not {(?n)#pragma omp target .* defaultmap\(firstprivate\) firstprivate\(_ZTI2C2\) firstprivate\(_ZTI2C1\) firstprivate\(_ZTV2C1\)$} gimple { xfail *-*-* } } } */ diff --git a/libgomp/testsuite/libgomp.c++/pr119692-1-5.C b/libgomp/testsuite/libgomp.c++/pr119692-1-5.C new file mode 100644 index 0000000..02121b6 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/pr119692-1-5.C @@ -0,0 +1,10 @@ +/* PR119692 "C++ 'typeinfo', 'vtable' vs. OpenACC, OpenMP 'target' offloading" */ + +/* { dg-additional-options -DDEFAULT=defaultmap(to) } + Wrong code for offloading execution. + { dg-xfail-run-if PR119692 { offload_device } } */ +/* { dg-additional-options -fdump-tree-gimple } */ + +#include "pr119692-1-1.C" + +/* { dg-final { scan-tree-dump-not {(?n)#pragma omp target .* defaultmap\(to\) map\(to:_ZTI2C2 \[len: [0-9]+\] \[runtime_implicit\]\) map\(to:_ZTI2C1 \[len: [0-9]+\] \[runtime_implicit\]\) map\(to:_ZTV2C1 \[len: [0-9]+\] \[runtime_implicit\]\)$} gimple { xfail *-*-* } } } */ diff --git a/libgomp/testsuite/libgomp.c++/pr96390.C b/libgomp/testsuite/libgomp.c++/pr96390.C index 1f3c3e0..be19601 100644 --- a/libgomp/testsuite/libgomp.c++/pr96390.C +++ b/libgomp/testsuite/libgomp.c++/pr96390.C @@ -1,6 +1,4 @@ /* { dg-additional-options "-O0 -fdump-tree-omplower" } */ -/* { dg-additional-options "-foffload=-Wa,--verify" { target offload_target_nvptx } } */ -/* { dg-xfail-if "PR 97106/PR 97102 - .alias not (yet) supported for nvptx" { offload_target_nvptx } } */ #include <cstdlib> #include <type_traits> diff --git a/libgomp/testsuite/libgomp.c++/target-exceptions-bad_cast-1.C b/libgomp/testsuite/libgomp.c++/target-exceptions-bad_cast-1.C new file mode 100644 index 0000000..3848295 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-exceptions-bad_cast-1.C @@ -0,0 +1,25 @@ +/* 'std::bad_cast' exception in OpenMP 'target' region. */ + +/* { dg-require-effective-target exceptions } + { dg-additional-options -fexceptions } */ +/* { dg-additional-options -fdump-tree-optimized-raw } + { dg-additional-options -foffload-options=-fdump-tree-optimized-raw } */ + +#include "../libgomp.oacc-c++/exceptions-bad_cast-1.C" + +/* { dg-output {CheCKpOInT[\r\n]+} } + + { dg-final { scan-tree-dump-times {gimple_call <__cxa_bad_cast, } 1 optimized } } + { dg-final { scan-offload-tree-dump-times {gimple_call <__cxa_bad_cast, } 1 optimized } } + For host execution, we print something like: + terminate called after throwing an instance of 'std::bad_cast' + what(): std::bad_cast + Aborted (core dumped) + { dg-output {.*std::bad_cast} { target { ! offload_device } } } + For GCN, nvptx offload execution, we don't print anything, but just 'abort'. + + TODO For GCN, nvptx offload execution, this currently doesn't 'abort' due to + the 'std::bad_cast' exception, but rather due to SIGSEGV in 'dynamic_cast'; + PR119692. + + { dg-shouldfail {'std::bad_cast' exception} } */ diff --git a/libgomp/testsuite/libgomp.c++/target-exceptions-bad_cast-2-offload-sorry-GCN.C b/libgomp/testsuite/libgomp.c++/target-exceptions-bad_cast-2-offload-sorry-GCN.C new file mode 100644 index 0000000..93884df --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-exceptions-bad_cast-2-offload-sorry-GCN.C @@ -0,0 +1,19 @@ +/* 'std::bad_cast' exception in OpenMP 'target' region, caught, '-foffload-options=-mno-fake-exceptions'. */ + +/* As this test case involves an expected offload compilation failure, we have to handle each offload target individually. + { dg-do link { target offload_target_amdgcn } } + { dg-additional-options -foffload=amdgcn-amdhsa } */ +/* { dg-require-effective-target exceptions } + { dg-additional-options -fexceptions } */ +/* { dg-additional-options -foffload-options=-mno-fake-exceptions } */ +/* { dg-additional-options -fdump-tree-optimized-raw } + { dg-additional-options -foffload-options=-fdump-tree-optimized-raw } */ + +#include "target-exceptions-bad_cast-2.C" + +/* { dg-final { scan-tree-dump-times {gimple_call <__cxa_bad_cast, } 1 optimized } } + { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump-times {gimple_call <__cxa_bad_cast, } 1 optimized } } + Given '-foffload-options=-mno-fake-exceptions', offload compilation fails: + { dg-regexp {[^\r\n]+: In function 'main[^']+':[\r\n]+(?:[^\r\n]+: sorry, unimplemented: exception handling not supported[\r\n]+)+} } + (Note, using 'dg-regexp' instead of 'dg-message', as the former runs before the auto-mark-UNSUPPORTED.) + { dg-excess-errors {'mkoffload' failure etc.} } */ diff --git a/libgomp/testsuite/libgomp.c++/target-exceptions-bad_cast-2-offload-sorry-nvptx.C b/libgomp/testsuite/libgomp.c++/target-exceptions-bad_cast-2-offload-sorry-nvptx.C new file mode 100644 index 0000000..83ec89b --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-exceptions-bad_cast-2-offload-sorry-nvptx.C @@ -0,0 +1,19 @@ +/* 'std::bad_cast' exception in OpenMP 'target' region, caught, '-foffload-options=-mno-fake-exceptions'. */ + +/* As this test case involves an expected offload compilation failure, we have to handle each offload target individually. + { dg-do link { target offload_target_nvptx } } + { dg-additional-options -foffload=nvptx-none } */ +/* { dg-require-effective-target exceptions } + { dg-additional-options -fexceptions } */ +/* { dg-additional-options -foffload-options=-mno-fake-exceptions } */ +/* { dg-additional-options -fdump-tree-optimized-raw } + { dg-additional-options -foffload-options=-fdump-tree-optimized-raw } */ + +#include "target-exceptions-bad_cast-2.C" + +/* { dg-final { scan-tree-dump-times {gimple_call <__cxa_bad_cast, } 1 optimized } } + { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump-times {gimple_call <__cxa_bad_cast, } 1 optimized } } + Given '-foffload-options=-mno-fake-exceptions', offload compilation fails: + { dg-regexp {[^\r\n]+: In function 'main[^']+':[\r\n]+(?:[^\r\n]+: sorry, unimplemented: exception handling not supported[\r\n]+)+} } + (Note, using 'dg-regexp' instead of 'dg-message', as the former runs before the auto-mark-UNSUPPORTED.) + { dg-excess-errors {'mkoffload' failure etc.} } */ diff --git a/libgomp/testsuite/libgomp.c++/target-exceptions-bad_cast-2.C b/libgomp/testsuite/libgomp.c++/target-exceptions-bad_cast-2.C new file mode 100644 index 0000000..8861740 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-exceptions-bad_cast-2.C @@ -0,0 +1,24 @@ +/* 'std::bad_cast' exception in OpenMP 'target' region, caught. */ + +/* { dg-require-effective-target exceptions } + { dg-additional-options -fexceptions } */ +/* { dg-additional-options -fdump-tree-optimized-raw } + { dg-additional-options -foffload-options=-fdump-tree-optimized-raw } */ +/* { dg-bogus {_ZTISt8bad_cast} PR119734 { target offload_target_nvptx xfail *-*-* } 0 } + { dg-excess-errors {'mkoffload' failure etc.} { xfail offload_target_nvptx } } */ + +#include "../libgomp.oacc-c++/exceptions-bad_cast-2.C" + +/* { dg-output {CheCKpOInT[\r\n]+} } + + { dg-final { scan-tree-dump-times {gimple_call <__cxa_bad_cast, } 1 optimized } } + { dg-final { scan-offload-tree-dump-times {gimple_call <__cxa_bad_cast, } 1 optimized } } + { dg-output {.*caught 'std::bad_cast'[\r\n]+} { target { ! offload_device } } } + For GCN, nvptx offload execution, we don't print anything, but just 'abort'. + + TODO For GCN, nvptx offload execution, this currently doesn't 'abort' due to + the 'std::bad_cast' exception, but rather due to SIGSEGV in 'dynamic_cast'; + PR119692. + + For GCN, nvptx offload execution, there is no 'catch'ing; any exception is fatal. + { dg-shouldfail {'MyException' exception} { offload_device } } */ diff --git a/libgomp/testsuite/libgomp.c++/target-exceptions-bad_cast-3.C b/libgomp/testsuite/libgomp.c++/target-exceptions-bad_cast-3.C new file mode 100644 index 0000000..efed64f --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-exceptions-bad_cast-3.C @@ -0,0 +1,17 @@ +/* 'std::bad_cast' exception in OpenMP 'target' region, dead code. */ + +/* { dg-require-effective-target exceptions } + { dg-additional-options -fexceptions } */ +/* { dg-additional-options -DDEFAULT=defaultmap(to) } + ... to avoid wrong code for offloading execution; PR119692. + With this, the device code still isn't correct, but the defects are in dead code. + { dg-additional-options -fdump-tree-gimple } */ +/* { dg-additional-options -fdump-tree-optimized-raw } + { dg-additional-options -foffload-options=-fdump-tree-optimized-raw } */ + +#include "../libgomp.oacc-c++/exceptions-bad_cast-3.C" + +/* { dg-final { scan-tree-dump-not {(?n)#pragma omp target .* defaultmap\(to\) map\(to:_ZTI2C2 \[len: [0-9]+\] \[runtime_implicit\]\) map\(to:_ZTI2C1 \[len: [0-9]+\] \[runtime_implicit\]\) map\(to:_ZTV2C1 \[len: [0-9]+\] \[runtime_implicit\]\)$} gimple { xfail *-*-* } } } */ + +/* { dg-final { scan-tree-dump-times {gimple_call <__cxa_bad_cast, } 1 optimized } } + { dg-final { scan-offload-tree-dump-times {gimple_call <__cxa_bad_cast, } 1 optimized } } */ diff --git a/libgomp/testsuite/libgomp.c++/target-exceptions-pr118794-1-offload-sorry-GCN.C b/libgomp/testsuite/libgomp.c++/target-exceptions-pr118794-1-offload-sorry-GCN.C new file mode 100644 index 0000000..3cdedf4 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-exceptions-pr118794-1-offload-sorry-GCN.C @@ -0,0 +1,24 @@ +/* Exception handling constructs in dead code, '-foffload-options=-mno-fake-exceptions'. */ + +/* As this test case involves an expected offload compilation failure, we have to handle each offload target individually. + { dg-do link { target offload_target_amdgcn } } + { dg-additional-options -foffload=amdgcn-amdhsa } */ +/* { dg-require-effective-target exceptions } + { dg-additional-options -fexceptions } */ +/* { dg-additional-options -foffload-options=-mno-fake-exceptions } */ +/* { dg-additional-options -O0 } */ +/* { dg-additional-options -fdump-tree-optimized-raw } + { dg-additional-options -foffload-options=-fdump-tree-optimized-raw } */ + +#include "target-exceptions-pr118794-1.C" + +/* In this specific C++ arrangement, distilled from PR118794, GCC synthesizes + '__builtin_eh_pointer', '__builtin_unwind_resume' calls as dead code in 'f': + { dg-final { scan-tree-dump-times {gimple_call <__builtin_eh_pointer, } 1 optimized } } + { dg-final { scan-tree-dump-times {gimple_call <__builtin_unwind_resume, } 1 optimized } } + { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump-times {gimple_call <__builtin_eh_pointer, } 1 optimized } } + { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump-times {gimple_call <__builtin_unwind_resume, } 1 optimized } } + Given '-O0' and '-foffload-options=-mno-fake-exceptions', offload compilation fails: + { dg-regexp {[^\r\n]+: In function 'f':[\r\n]+(?:[^\r\n]+: sorry, unimplemented: exception handling not supported[\r\n]+)+} } + (Note, using 'dg-regexp' instead of 'dg-message', as the former runs before the auto-mark-UNSUPPORTED.) + { dg-excess-errors {'mkoffload' failure etc.} } */ diff --git a/libgomp/testsuite/libgomp.c++/target-exceptions-pr118794-1-offload-sorry-nvptx.C b/libgomp/testsuite/libgomp.c++/target-exceptions-pr118794-1-offload-sorry-nvptx.C new file mode 100644 index 0000000..ef996cf --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-exceptions-pr118794-1-offload-sorry-nvptx.C @@ -0,0 +1,24 @@ +/* Exception handling constructs in dead code, '-foffload-options=-mno-fake-exceptions'. */ + +/* As this test case involves an expected offload compilation failure, we have to handle each offload target individually. + { dg-do link { target offload_target_nvptx } } + { dg-additional-options -foffload=nvptx-none } */ +/* { dg-require-effective-target exceptions } + { dg-additional-options -fexceptions } */ +/* { dg-additional-options -foffload-options=-mno-fake-exceptions } */ +/* { dg-additional-options -O0 } */ +/* { dg-additional-options -fdump-tree-optimized-raw } + { dg-additional-options -foffload-options=-fdump-tree-optimized-raw } */ + +#include "target-exceptions-pr118794-1.C" + +/* In this specific C++ arrangement, distilled from PR118794, GCC synthesizes + '__builtin_eh_pointer', '__builtin_unwind_resume' calls as dead code in 'f': + { dg-final { scan-tree-dump-times {gimple_call <__builtin_eh_pointer, } 1 optimized } } + { dg-final { scan-tree-dump-times {gimple_call <__builtin_unwind_resume, } 1 optimized } } + { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump-times {gimple_call <__builtin_eh_pointer, } 1 optimized } } + { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump-times {gimple_call <__builtin_unwind_resume, } 1 optimized } } + Given '-O0' and '-foffload-options=-mno-fake-exceptions', offload compilation fails: + { dg-regexp {[^\r\n]+: In function 'f':[\r\n]+(?:[^\r\n]+: sorry, unimplemented: exception handling not supported[\r\n]+)+} } + (Note, using 'dg-regexp' instead of 'dg-message', as the former runs before the auto-mark-UNSUPPORTED.) + { dg-excess-errors {'mkoffload' failure etc.} } */ diff --git a/libgomp/testsuite/libgomp.c++/target-exceptions-pr118794-1.C b/libgomp/testsuite/libgomp.c++/target-exceptions-pr118794-1.C new file mode 100644 index 0000000..24e3d07 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-exceptions-pr118794-1.C @@ -0,0 +1,57 @@ +/* Exception handling constructs in dead code. */ + +/* { dg-require-effective-target exceptions } + { dg-additional-options -fexceptions } */ +/* { dg-additional-options -O0 } */ +/* { dg-additional-options -fdump-tree-optimized-raw } + { dg-additional-options -foffload-options=-fdump-tree-optimized-raw } */ + +/* See also '../../../gcc/testsuite/g++.target/gcn/exceptions-pr118794-1.C', + '../../../gcc/testsuite/g++.target/nvptx/exceptions-pr118794-1.C'. */ + +#pragma omp begin declare target + +bool ok = false; + +template <typename T> +struct C +{ + C() + { + ok = true; + } + C(int) {}; + ~C() {}; + + __attribute__((noipa)) + void m() + { + C c; + } +}; + +inline void f() +{ + C<double> c(1); + c.m(); +} + +#pragma omp end declare target + +int main() +{ +#pragma omp target + { + f(); + } +#pragma omp target update from(ok) + if (!ok) + __builtin_abort(); +} + +/* In this specific C++ arrangement, distilled from PR118794, GCC synthesizes + '__builtin_eh_pointer', '__builtin_unwind_resume' calls as dead code in 'f': + { dg-final { scan-tree-dump-times {gimple_call <__builtin_eh_pointer, } 1 optimized } } + { dg-final { scan-tree-dump-times {gimple_call <__builtin_unwind_resume, } 1 optimized } } + { dg-final { scan-offload-tree-dump-times {gimple_call <__builtin_eh_pointer, } 1 optimized } } + { dg-final { scan-offload-tree-dump-times {gimple_call <__builtin_unwind_resume, } 1 optimized } } */ diff --git a/libgomp/testsuite/libgomp.c++/target-exceptions-throw-1-O0.C b/libgomp/testsuite/libgomp.c++/target-exceptions-throw-1-O0.C new file mode 100644 index 0000000..00d7c13 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-exceptions-throw-1-O0.C @@ -0,0 +1,23 @@ +/* 'throw' in OpenMP 'target' region. */ + +/* { dg-additional-options -O0 } */ +/* { dg-require-effective-target exceptions } + { dg-additional-options -fexceptions } */ +/* { dg-additional-options -fdump-tree-optimized-raw } + { dg-additional-options -foffload-options=-fdump-tree-optimized-raw } */ + +#include "target-exceptions-throw-1.C" + +/* { dg-output {CheCKpOInT[\r\n]+} } + + { dg-final { scan-tree-dump-times {gimple_call <__cxa_allocate_exception, } 1 optimized } } + { dg-final { scan-tree-dump-times {gimple_call <__cxa_throw, } 1 optimized } } + { dg-final { scan-offload-tree-dump-times {gimple_call <__cxa_allocate_exception, } 1 optimized } } + { dg-final { scan-offload-tree-dump-times {gimple_call <__cxa_throw, } 1 optimized } } + For host execution, we print something like: + terminate called after throwing an instance of 'MyException' + Aborted (core dumped) + { dg-output {.*MyException} { target { ! offload_device } } } + For GCN, nvptx offload execution, we don't print anything, but just 'abort'. + + { dg-shouldfail {'MyException' exception} } */ diff --git a/libgomp/testsuite/libgomp.c++/target-exceptions-throw-1.C b/libgomp/testsuite/libgomp.c++/target-exceptions-throw-1.C new file mode 100644 index 0000000..2467061 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-exceptions-throw-1.C @@ -0,0 +1,25 @@ +/* 'throw' in OpenMP 'target' region. */ + +/* { dg-require-effective-target exceptions } + { dg-additional-options -fexceptions } */ +/* { dg-additional-options -fdump-tree-optimized-raw } + { dg-additional-options -foffload-options=-fdump-tree-optimized-raw } */ +/* { dg-bogus {Size expression must be absolute\.} PR119737 { target offload_target_amdgcn xfail *-*-* } 0 } + { dg-ice PR119737 { offload_target_amdgcn } } + { dg-excess-errors {'mkoffload' failures etc.} { xfail offload_target_amdgcn } } */ + +#include "../libgomp.oacc-c++/exceptions-throw-1.C" + +/* { dg-output {CheCKpOInT[\r\n]+} } + + { dg-final { scan-tree-dump-times {gimple_call <__cxa_allocate_exception, } 1 optimized } } + { dg-final { scan-tree-dump-times {gimple_call <__cxa_throw, } 1 optimized } } + { dg-final { scan-offload-tree-dump-times {gimple_call <__cxa_allocate_exception, } 1 optimized } } + { dg-final { scan-offload-tree-dump-times {gimple_call <__cxa_throw, } 1 optimized } } + For host execution, we print something like: + terminate called after throwing an instance of 'MyException' + Aborted (core dumped) + { dg-output {.*MyException} { target { ! offload_device } } } + For GCN, nvptx offload execution, we don't print anything, but just 'abort'. + + { dg-shouldfail {'MyException' exception} } */ diff --git a/libgomp/testsuite/libgomp.c++/target-exceptions-throw-2-O0.C b/libgomp/testsuite/libgomp.c++/target-exceptions-throw-2-O0.C new file mode 100644 index 0000000..b7a311d --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-exceptions-throw-2-O0.C @@ -0,0 +1,25 @@ +/* 'throw' in OpenMP 'target' region, caught. */ + +/* { dg-additional-options -O0 } */ +/* { dg-require-effective-target exceptions } + { dg-additional-options -fexceptions } */ +/* { dg-additional-options -fdump-tree-optimized-raw } + { dg-additional-options -foffload-options=-fdump-tree-optimized-raw } */ +/* { dg-bogus {undefined symbol: typeinfo name for MyException} PR119806 { target offload_target_amdgcn xfail *-*-* } 0 } + { dg-excess-errors {'mkoffload' failure etc.} { xfail offload_target_amdgcn } } */ +/* { dg-bogus {Initial value type mismatch} PR119806 { target offload_target_nvptx xfail *-*-* } 0 } + { dg-excess-errors {'mkoffload' failure etc.} { xfail offload_target_nvptx } } */ + +#include "target-exceptions-throw-2.C" + +/* { dg-output {CheCKpOInT[\r\n]+} } + + { dg-final { scan-tree-dump-times {gimple_call <__cxa_allocate_exception, } 1 optimized } } + { dg-final { scan-tree-dump-times {gimple_call <__cxa_throw, } 1 optimized } } + { dg-final { scan-offload-tree-dump-times {gimple_call <__cxa_allocate_exception, } 1 optimized } } + { dg-final { scan-offload-tree-dump-times {gimple_call <__cxa_throw, } 1 optimized } } + { dg-output {.*caught 'MyException'[\r\n]+} { target { ! offload_device } } } + For GCN, nvptx offload execution, we don't print anything, but just 'abort'. + + For GCN, nvptx offload execution, there is no 'catch'ing; any exception is fatal. + { dg-shouldfail {'MyException' exception} { offload_device } } */ diff --git a/libgomp/testsuite/libgomp.c++/target-exceptions-throw-2-offload-sorry-GCN.C b/libgomp/testsuite/libgomp.c++/target-exceptions-throw-2-offload-sorry-GCN.C new file mode 100644 index 0000000..9905b1f --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-exceptions-throw-2-offload-sorry-GCN.C @@ -0,0 +1,21 @@ +/* 'throw' in OpenMP 'target' region, caught, -foffload-options=-mno-fake-exceptions. */ + +/* As this test case involves an expected offload compilation failure, we have to handle each offload target individually. + { dg-do link { target offload_target_amdgcn } } + { dg-additional-options -foffload=amdgcn-amdhsa } */ +/* { dg-require-effective-target exceptions } + { dg-additional-options -fexceptions } */ +/* { dg-additional-options -foffload-options=-mno-fake-exceptions } */ +/* { dg-additional-options -fdump-tree-optimized-raw } + { dg-additional-options -foffload-options=-fdump-tree-optimized-raw } */ + +#include "target-exceptions-throw-2.C" + +/* { dg-final { scan-tree-dump-times {gimple_call <__cxa_allocate_exception, } 1 optimized } } + { dg-final { scan-tree-dump-times {gimple_call <__cxa_throw, } 1 optimized } } + { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump-times {gimple_call <__cxa_allocate_exception, } 1 optimized } } + { dg-final { only_for_offload_target amdgcn-amdhsa scan-offload-tree-dump-times {gimple_call <__cxa_throw, } 1 optimized } } + Given '-foffload-options=-mno-fake-exceptions', offload compilation fails: + { dg-regexp {[^\r\n]+: In function 'main[^']+':[\r\n]+(?:[^\r\n]+: sorry, unimplemented: exception handling not supported[\r\n]+)+} } + (Note, using 'dg-regexp' instead of 'dg-message', as the former runs before the auto-mark-UNSUPPORTED.) + { dg-excess-errors {'mkoffload' failure etc.} } */ diff --git a/libgomp/testsuite/libgomp.c++/target-exceptions-throw-2-offload-sorry-nvptx.C b/libgomp/testsuite/libgomp.c++/target-exceptions-throw-2-offload-sorry-nvptx.C new file mode 100644 index 0000000..da267d6 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-exceptions-throw-2-offload-sorry-nvptx.C @@ -0,0 +1,21 @@ +/* 'throw' in OpenMP 'target' region, caught, '-foffload-options=-mno-fake-exceptions'. */ + +/* As this test case involves an expected offload compilation failure, we have to handle each offload target individually. + { dg-do link { target offload_target_nvptx } } + { dg-additional-options -foffload=nvptx-none } */ +/* { dg-require-effective-target exceptions } + { dg-additional-options -fexceptions } */ +/* { dg-additional-options -foffload-options=-mno-fake-exceptions } */ +/* { dg-additional-options -fdump-tree-optimized-raw } + { dg-additional-options -foffload-options=-fdump-tree-optimized-raw } */ + +#include "target-exceptions-throw-2.C" + +/* { dg-final { scan-tree-dump-times {gimple_call <__cxa_allocate_exception, } 1 optimized } } + { dg-final { scan-tree-dump-times {gimple_call <__cxa_throw, } 1 optimized } } + { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump-times {gimple_call <__cxa_allocate_exception, } 1 optimized } } + { dg-final { only_for_offload_target nvptx-none scan-offload-tree-dump-times {gimple_call <__cxa_throw, } 1 optimized } } + Given '-foffload-options=-mno-fake-exceptions', offload compilation fails: + { dg-regexp {[^\r\n]+: In function 'main[^']+':[\r\n]+(?:[^\r\n]+: sorry, unimplemented: exception handling not supported[\r\n]+)+} } + (Note, using 'dg-regexp' instead of 'dg-message', as the former runs before the auto-mark-UNSUPPORTED.) + { dg-excess-errors {'mkoffload' failure etc.} } */ diff --git a/libgomp/testsuite/libgomp.c++/target-exceptions-throw-2.C b/libgomp/testsuite/libgomp.c++/target-exceptions-throw-2.C new file mode 100644 index 0000000..e85e6c3 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-exceptions-throw-2.C @@ -0,0 +1,23 @@ +/* 'throw' in OpenMP 'target' region, caught. */ + +/* { dg-require-effective-target exceptions } + { dg-additional-options -fexceptions } */ +/* { dg-additional-options -fdump-tree-optimized-raw } + { dg-additional-options -foffload-options=-fdump-tree-optimized-raw } */ +/* { dg-bogus {Size expression must be absolute\.} PR119737 { target offload_target_amdgcn xfail *-*-* } 0 } + { dg-ice PR119737 { offload_target_amdgcn } } + { dg-excess-errors {'mkoffload' failures etc.} { xfail offload_target_amdgcn } } */ + +#include "../libgomp.oacc-c++/exceptions-throw-2.C" + +/* { dg-output {CheCKpOInT[\r\n]+} } + + { dg-final { scan-tree-dump-times {gimple_call <__cxa_allocate_exception, } 1 optimized } } + { dg-final { scan-tree-dump-times {gimple_call <__cxa_throw, } 1 optimized } } + { dg-final { scan-offload-tree-dump-times {gimple_call <__cxa_allocate_exception, } 1 optimized } } + { dg-final { scan-offload-tree-dump-times {gimple_call <__cxa_throw, } 1 optimized } } + { dg-output {.*caught 'MyException'[\r\n]+} { target { ! offload_device } } } + For GCN, nvptx offload execution, we don't print anything, but just 'abort'. + + For GCN, nvptx offload execution, there is no 'catch'ing; any exception is fatal. + { dg-shouldfail {'MyException' exception} { offload_device } } */ diff --git a/libgomp/testsuite/libgomp.c++/target-exceptions-throw-3.C b/libgomp/testsuite/libgomp.c++/target-exceptions-throw-3.C new file mode 100644 index 0000000..c35180d --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-exceptions-throw-3.C @@ -0,0 +1,19 @@ +/* 'throw' in OpenMP 'target' region, dead code. */ + +/* { dg-require-effective-target exceptions } + { dg-additional-options -fexceptions } */ +/* { dg-additional-options -DDEFAULT=defaultmap(to) } + ... to avoid wrong code for offloading execution; PR119692. + With this, the device code still isn't correct, but the defects are in dead code. + { dg-additional-options -fdump-tree-gimple } */ +/* { dg-additional-options -fdump-tree-optimized-raw } + { dg-additional-options -foffload-options=-fdump-tree-optimized-raw } */ + +#include "../libgomp.oacc-c++/exceptions-throw-3.C" + +/* { dg-final { scan-tree-dump-not {(?n)#pragma omp target .* defaultmap\(to\) map\(to:_ZTI11MyException \[len: [0-9]+\] \[runtime_implicit\]\)$} gimple { xfail *-*-* } } } */ + +/* { dg-final { scan-tree-dump-times {gimple_call <__cxa_allocate_exception, } 1 optimized } } + { dg-final { scan-tree-dump-times {gimple_call <__cxa_throw, } 1 optimized } } + { dg-final { scan-offload-tree-dump-times {gimple_call <__cxa_allocate_exception, } 1 optimized } } + { dg-final { scan-offload-tree-dump-times {gimple_call <__cxa_throw, } 1 optimized } } */ |