diff options
Diffstat (limited to 'libgomp')
51 files changed, 4189 insertions, 6 deletions
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 9d9ecfb..096e17b 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,157 @@ +2025-04-17 Jakub Jelinek <jakub@redhat.com> + + PR libgomp/119849 + * testsuite/libgomp.c++/allocator-1.C (test_inequality, main): Guard + ompx::allocator::gnu_pinned_mem uses with #ifdef __gnu_linux__. + * testsuite/libgomp.c++/allocator-2.C (main): Likewise. + +2025-04-17 Tobias Burnus <tburnus@baylibre.com> + + * libgomp.texi (gcn interop, nvptx interop): For HIP with C/C++, add + a note about setting a preprocessor define. + +2025-04-16 Thomas Schwinge <tschwinge@baylibre.com> + + * testsuite/libgomp.c++/target-exceptions-pr118794-1.C: Remove + 'ALWAYS_INLINE' workaround. + +2025-04-16 Thomas Schwinge <tschwinge@baylibre.com> + + PR target/106445 + * testsuite/libgomp.c++/pr106445-1.C: New. + * testsuite/libgomp.c++/pr106445-1-O0.C: Likewise. + +2025-04-16 Thomas Schwinge <tschwinge@baylibre.com> + + PR target/97106 + * testsuite/libgomp.c++/pr96390.C: Un-XFAIL nvptx offloading. + * testsuite/libgomp.c-c++-common/pr96390.c: Adjust. + +2025-04-15 Tobias Burnus <tburnus@baylibre.com> + + * libgomp.texi (gcn, nvptx): Mention self_maps clause + besides unified_shared_memory in the requirements item. + +2025-04-15 waffl3x <waffl3x@baylibre.com> + + * omp.h.in: Add omp::allocator::* and ompx::allocator::* allocators. + (__detail::__allocator_templ<T, omp_allocator_handle_t>): + New struct template. + (null_allocator<T>): New struct template. + (default_mem<T>): Likewise. + (large_cap_mem<T>): Likewise. + (const_mem<T>): Likewise. + (high_bw_mem<T>): Likewise. + (low_lat_mem<T>): Likewise. + (cgroup_mem<T>): Likewise. + (pteam_mem<T>): Likewise. + (thread_mem<T>): Likewise. + (ompx::allocator::gnu_pinned_mem<T>): Likewise. + * testsuite/libgomp.c++/allocator-1.C: New test. + * testsuite/libgomp.c++/allocator-2.C: New test. + +2025-04-15 Tobias Burnus <tburnus@baylibre.com> + + * libgomp.texi (5.0 Impl. Status): Mark mapping alloc comps as 'Y'. + * testsuite/libgomp.fortran/allocatable-comp.f90: New test. + * testsuite/libgomp.fortran/map-alloc-comp-3.f90: New test. + * testsuite/libgomp.fortran/map-alloc-comp-4.f90: New test. + * testsuite/libgomp.fortran/map-alloc-comp-5.f90: New test. + * testsuite/libgomp.fortran/map-alloc-comp-6.f90: New test. + * testsuite/libgomp.fortran/map-alloc-comp-7.f90: New test. + * testsuite/libgomp.fortran/map-alloc-comp-8.f90: New test. + * testsuite/libgomp.fortran/map-alloc-comp-9.f90: New test. + +2025-04-14 Thomas Schwinge <tschwinge@baylibre.com> + + PR target/118794 + * testsuite/libgomp.c++/target-exceptions-bad_cast-2-offload-sorry-GCN.C: + Set '-foffload-options=-mno-fake-exceptions'. + * testsuite/libgomp.c++/target-exceptions-bad_cast-2-offload-sorry-nvptx.C: + Likewise. + * testsuite/libgomp.c++/target-exceptions-pr118794-1-offload-sorry-GCN.C: + Likewise. + * testsuite/libgomp.c++/target-exceptions-pr118794-1-offload-sorry-nvptx.C: + Likewise. + * testsuite/libgomp.c++/target-exceptions-throw-2-offload-sorry-GCN.C: + Likewise. + * testsuite/libgomp.c++/target-exceptions-throw-2-offload-sorry-nvptx.C: + Likewise. + * testsuite/libgomp.oacc-c++/exceptions-bad_cast-2-offload-sorry-GCN.C: + Likewise. + * testsuite/libgomp.oacc-c++/exceptions-bad_cast-2-offload-sorry-nvptx.C: + Likewise. + * testsuite/libgomp.oacc-c++/exceptions-throw-2-offload-sorry-GCN.C: + Likewise. + * testsuite/libgomp.oacc-c++/exceptions-throw-2-offload-sorry-nvptx.C: + Likewise. + * testsuite/libgomp.c++/target-exceptions-bad_cast-2.C: Adjust. + * testsuite/libgomp.c++/target-exceptions-pr118794-1.C: Likewise. + * testsuite/libgomp.c++/target-exceptions-throw-2.C: Likewise. + * testsuite/libgomp.oacc-c++/exceptions-bad_cast-2.C: Likewise. + * testsuite/libgomp.oacc-c++/exceptions-throw-2.C: Likewise. + * testsuite/libgomp.c++/target-exceptions-throw-2-O0.C: New. + +2025-04-14 Thomas Schwinge <tschwinge@baylibre.com> + + * testsuite/libgomp.c++/target-exceptions-throw-3.C: New. + * testsuite/libgomp.oacc-c++/exceptions-throw-3.C: Likewise. + +2025-04-14 Thomas Schwinge <tschwinge@baylibre.com> + + * testsuite/libgomp.c++/target-exceptions-throw-2.C: New. + * testsuite/libgomp.c++/target-exceptions-throw-2-offload-sorry-GCN.C: Likewise. + * testsuite/libgomp.c++/target-exceptions-throw-2-offload-sorry-nvptx.C: Likewise. + * testsuite/libgomp.oacc-c++/exceptions-throw-2.C: Likewise. + * testsuite/libgomp.oacc-c++/exceptions-throw-2-offload-sorry-GCN.C: Likewise. + * testsuite/libgomp.oacc-c++/exceptions-throw-2-offload-sorry-nvptx.C: Likewise. + +2025-04-14 Thomas Schwinge <tschwinge@baylibre.com> + + * testsuite/libgomp.c++/target-exceptions-throw-1.C: New. + * testsuite/libgomp.c++/target-exceptions-throw-1-O0.C: Likewise. + * testsuite/libgomp.oacc-c++/exceptions-throw-1.C: Likewise. + +2025-04-14 Thomas Schwinge <tschwinge@baylibre.com> + + * testsuite/libgomp.c++/target-exceptions-bad_cast-3.C: New. + * testsuite/libgomp.oacc-c++/exceptions-bad_cast-3.C: Likewise. + +2025-04-14 Thomas Schwinge <tschwinge@baylibre.com> + + * testsuite/libgomp.c++/target-exceptions-bad_cast-2.C: New. + * testsuite/libgomp.c++/target-exceptions-bad_cast-2-offload-sorry-GCN.C: Likewise. + * testsuite/libgomp.c++/target-exceptions-bad_cast-2-offload-sorry-nvptx.C: Likewise. + * testsuite/libgomp.oacc-c++/exceptions-bad_cast-2.C: Likewise. + * testsuite/libgomp.oacc-c++/exceptions-bad_cast-2-offload-sorry-GCN.C: Likewise. + * testsuite/libgomp.oacc-c++/exceptions-bad_cast-2-offload-sorry-nvptx.C: Likewise. + +2025-04-14 Thomas Schwinge <tschwinge@baylibre.com> + + * testsuite/libgomp.c++/target-exceptions-bad_cast-1.C: New. + * testsuite/libgomp.oacc-c++/exceptions-bad_cast-1.C: Likewise. + +2025-04-14 Thomas Schwinge <tschwinge@baylibre.com> + + PR target/118794 + * testsuite/libgomp.c++/target-exceptions-pr118794-1.C: New. + * testsuite/libgomp.c++/target-exceptions-pr118794-1-offload-sorry-GCN.C: + Likewise. + * testsuite/libgomp.c++/target-exceptions-pr118794-1-offload-sorry-nvptx.C: + Likewise. + +2025-04-14 Thomas Schwinge <tschwinge@baylibre.com> + + PR c++/119692 + * testsuite/libgomp.c++/pr119692-1-1.C: New. + * testsuite/libgomp.c++/pr119692-1-2.C: Likewise. + * testsuite/libgomp.c++/pr119692-1-3.C: Likewise. + * testsuite/libgomp.c++/pr119692-1-4.C: Likewise. + * testsuite/libgomp.c++/pr119692-1-5.C: Likewise. + * testsuite/libgomp.oacc-c++/pr119692-1-1.C: Likewise. + * testsuite/libgomp.oacc-c++/pr119692-1-2.C: Likewise. + * testsuite/libgomp.oacc-c++/pr119692-1-3.C: Likewise. + 2025-04-10 Richard Sandiford <richard.sandiford@arm.com> * testsuite/libgomp.c-target/aarch64/firstprivate.c: Add +sve pragma. diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi index fed9d5e..6909c2b 100644 --- a/libgomp/libgomp.texi +++ b/libgomp/libgomp.texi @@ -258,7 +258,7 @@ The OpenMP 4.5 specification is fully supported. device memory mapped by an array section @tab P @tab @item Mapping of Fortran pointer and allocatable variables, including pointer and allocatable components of variables - @tab P @tab Mapping of vars with allocatable components unsupported + @tab Y @tab @item @code{defaultmap} extensions @tab Y @tab @item @code{declare mapper} directive @tab N @tab @item @code{omp_get_supported_active_levels} routine @tab Y @tab @@ -6888,7 +6888,7 @@ The implementation remark: @code{device(ancestor:1)}) are processed serially per @code{target} region such that the next reverse offload region is only executed after the previous one returned. -@item OpenMP code that has a @code{requires} directive with +@item OpenMP code that has a @code{requires} directive with @code{self_maps} or @code{unified_shared_memory} is only supported if all AMD GPUs have the @code{HSA_AMD_SYSTEM_INFO_SVM_ACCESSIBLE_BY_DEFAULT} property; for discrete GPUs, this may require setting the @code{HSA_XNACK} environment @@ -6945,6 +6945,9 @@ or string (str) data type, call @code{omp_get_interop_int}, Note that @code{device_num} is the OpenMP device number while @code{device} is the HIP device number or HSA device handle. +When using HIP with C and C++, the @code{__HIP_PLATFORM_AMD__} preprocessor +macro must be defined before including the HIP header files. + For the API routine call, add the prefix @code{omp_ipr_} to the property name; for instance: @smallexample @@ -7045,7 +7048,7 @@ The implementation remark: Per device, reverse offload regions are processed serially such that the next reverse offload region is only executed after the previous one returned. -@item OpenMP code that has a @code{requires} directive with +@item OpenMP code that has a @code{requires} directive with @code{self_maps} or @code{unified_shared_memory} runs on nvptx devices if and only if all of those support the @code{pageableMemoryAccess} property;@footnote{ @uref{https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#um-requirements}} @@ -7107,6 +7110,9 @@ or string (str) data type, call @code{omp_get_interop_int}, Note that @code{device_num} is the OpenMP device number while @code{device} is the CUDA, CUDA Driver, or HIP device number. +When using HIP with C and C++, the @code{__HIP_PLATFORM_NVIDIA__} preprocessor +macro must be defined before including the HIP header files. + For the API routine call, add the prefix @code{omp_ipr_} to the property name; for instance: @smallexample diff --git a/libgomp/omp.h.in b/libgomp/omp.h.in index d5e8be4..8d17db1 100644 --- a/libgomp/omp.h.in +++ b/libgomp/omp.h.in @@ -432,4 +432,136 @@ extern const char *omp_get_uid_from_device (int) __GOMP_NOTHROW; } #endif +#if __cplusplus >= 201103L + +/* std::__throw_bad_alloc and std::__throw_bad_array_new_length. */ +#include <bits/functexcept.h> + +namespace omp +{ +namespace allocator +{ + +namespace __detail +{ + +template<typename __T, omp_allocator_handle_t __Handle> +struct __allocator_templ +{ + using value_type = __T; + using pointer = __T*; + using const_pointer = const __T*; + using size_type = __SIZE_TYPE__; + using difference_type = __PTRDIFF_TYPE__; + + __T* + allocate (size_type __n) + { + if (__SIZE_MAX__ / sizeof(__T) < __n) + std::__throw_bad_array_new_length (); + void *__p = omp_aligned_alloc (alignof(__T), __n * sizeof(__T), __Handle); + if (!__p) + std::__throw_bad_alloc (); + return static_cast<__T*>(__p); + } + + void + deallocate (__T *__p, size_type) __GOMP_NOTHROW + { + omp_free (static_cast<void*>(__p), __Handle); + } +}; + +template<typename __T, typename __U, omp_allocator_handle_t __Handle> +constexpr bool +operator== (const __allocator_templ<__T, __Handle>&, + const __allocator_templ<__U, __Handle>&) __GOMP_NOTHROW +{ + return true; +} + +template<typename __T, omp_allocator_handle_t __Handle, + typename __U, omp_allocator_handle_t __UHandle> +constexpr bool +operator== (const __allocator_templ<__T, __Handle>&, + const __allocator_templ<__U, __UHandle>&) __GOMP_NOTHROW +{ + return false; +} + +template<typename __T, typename __U, omp_allocator_handle_t __Handle> +constexpr bool +operator!= (const __allocator_templ<__T, __Handle>&, + const __allocator_templ<__U, __Handle>&) __GOMP_NOTHROW +{ + return false; +} + +template<typename __T, omp_allocator_handle_t __Handle, + typename __U, omp_allocator_handle_t __UHandle> +constexpr bool +operator!= (const __allocator_templ<__T, __Handle>&, + const __allocator_templ<__U, __UHandle>&) __GOMP_NOTHROW +{ + return true; +} + +} /* namespace __detail */ + +template<typename __T> +struct null_allocator + : __detail::__allocator_templ<__T, omp_null_allocator> {}; + +template<typename __T> +struct default_mem + : __detail::__allocator_templ<__T, omp_default_mem_alloc> {}; + +template<typename __T> +struct large_cap_mem + : __detail::__allocator_templ<__T, omp_large_cap_mem_alloc> {}; + +template<typename __T> +struct const_mem + : __detail::__allocator_templ<__T, omp_const_mem_alloc> {}; + +template<typename __T> +struct high_bw_mem + : __detail::__allocator_templ<__T, omp_high_bw_mem_alloc> {}; + +template<typename __T> +struct low_lat_mem + : __detail::__allocator_templ<__T, omp_low_lat_mem_alloc> {}; + +template<typename __T> +struct cgroup_mem + : __detail::__allocator_templ<__T, omp_cgroup_mem_alloc> {}; + +template<typename __T> +struct pteam_mem + : __detail::__allocator_templ<__T, omp_pteam_mem_alloc> {}; + +template<typename __T> +struct thread_mem + : __detail::__allocator_templ<__T, omp_thread_mem_alloc> {}; + +} /* namespace allocator */ + +} /* namespace omp */ + +namespace ompx +{ + +namespace allocator +{ + +template<typename __T> +struct gnu_pinned_mem + : omp::allocator::__detail::__allocator_templ<__T, ompx_gnu_pinned_mem_alloc> {}; + +} /* namespace allocator */ + +} /* namespace ompx */ + +#endif /* __cplusplus */ + #endif /* _OMP_H */ 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 } } */ diff --git a/libgomp/testsuite/libgomp.c-c++-common/pr96390.c b/libgomp/testsuite/libgomp.c-c++-common/pr96390.c index b89f934..ca7865d 100644 --- a/libgomp/testsuite/libgomp.c-c++-common/pr96390.c +++ b/libgomp/testsuite/libgomp.c-c++-common/pr96390.c @@ -1,7 +1,7 @@ /* { dg-additional-options "-O0 -fdump-tree-omplower" } */ /* { dg-additional-options "-foffload=-Wa,--verify" { target offload_target_nvptx } } */ /* { dg-require-alias "" } */ -/* { dg-xfail-if "PR 97102/PR 97106 - .alias not (yet) supported for nvptx" { offload_target_nvptx } } */ +/* { dg-xfail-if PR105018 { offload_target_nvptx } } */ #ifdef __cplusplus extern "C" { diff --git a/libgomp/testsuite/libgomp.fortran/allocatable-comp.f90 b/libgomp/testsuite/libgomp.fortran/allocatable-comp.f90 new file mode 100644 index 0000000..383ecba --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/allocatable-comp.f90 @@ -0,0 +1,53 @@ +implicit none +type t + integer, allocatable :: a, b(:) +end type t +type(t) :: x, y, z +integer :: i + +!$omp target map(to: x) + if (allocated(x%a)) stop 1 + if (allocated(x%b)) stop 2 +!$omp end target + +allocate(x%a, x%b(-4:6)) +x%b(:) = [(i, i=-4,6)] + +!$omp target map(to: x) + if (.not. allocated(x%a)) stop 3 + if (.not. allocated(x%b)) stop 4 + if (lbound(x%b,1) /= -4) stop 5 + if (ubound(x%b,1) /= 6) stop 6 + if (any (x%b /= [(i, i=-4,6)])) stop 7 +!$omp end target + + +! The following only works with arrays due to +! PR fortran/96668 + +!$omp target enter data map(to: y, z) + +!$omp target map(to: y, z) + if (allocated(y%b)) stop 8 + if (allocated(z%b)) stop 9 +!$omp end target + +allocate(y%b(5), z%b(3)) +y%b = 42 +z%b = 99 + +! (implicitly) 'tofrom' mapped +! Planned for OpenMP 6.0 (but common extension) +! OpenMP <= 5.0 unclear +!$omp target map(to: y) + if (.not.allocated(y%b)) stop 10 + if (any (y%b /= 42)) stop 11 +!$omp end target + +! always map: OpenMP 5.1 (clarified) +!$omp target map(always, tofrom: z) + if (.not.allocated(z%b)) stop 12 + if (any (z%b /= 99)) stop 13 +!$omp end target + +end diff --git a/libgomp/testsuite/libgomp.fortran/map-alloc-comp-3.f90 b/libgomp/testsuite/libgomp.fortran/map-alloc-comp-3.f90 new file mode 100644 index 0000000..9d48c7c --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/map-alloc-comp-3.f90 @@ -0,0 +1,121 @@ +type t2 + integer x, y, z +end type t2 +type t + integer, allocatable :: A + integer, allocatable :: B(:) + type(t2), allocatable :: C + type(t2), allocatable :: D(:,:) +end type t + +type t3 + type(t) :: Q + type(t) :: R(5) +end type + +type(t) :: var, var2 +type(t3) :: var3, var4 + +! -------------------------------------- +! Assign + allocate +var%A = 45 +var%B = [1,2,3] +var%C = t2(6,5,4) +var%D = reshape([t2(1,2,3), t2(4,5,6), t2(11,12,13), t2(14,15,16)], [2,2]) + +! Assign + allocate +var2%A = 145 +var2%B = [991,992,993] +var2%C = t2(996,995,994) +var2%D = reshape([t2(199,299,399), t2(499,599,699), t2(1199,1299,1399), t2(1499,1599,1699)], [2,2]) + + +!$omp target map(to: var) map(tofrom: var2) + call foo(var, var2) +!$omp end target + +if (var2%A /= 45) stop 9 +if (any (var2%B /= [1,2,3])) stop 10 +if (var2%C%x /= 6) stop 11 +if (var2%C%y /= 5) stop 11 +if (var2%C%z /= 4) stop 11 +if (any (var2%D(:,:)%x /= reshape([1, 4, 11, 14], [2,2]))) stop 12 +if (any (var2%D(:,:)%y /= reshape([2, 5, 12, 15], [2,2]))) stop 12 +if (any (var2%D(:,:)%z /= reshape([3, 6, 13, 16], [2,2]))) stop 12 + +! -------------------------------------- +! Assign + allocate +var3%Q%A = 45 +var3%Q%B = [1,2,3] +var3%Q%C = t2(6,5,4) +var3%Q%D = reshape([t2(1,2,3), t2(4,5,6), t2(11,12,13), t2(14,15,16)], [2,2]) + +var3%R(2)%A = 45 +var3%R(2)%B = [1,2,3] +var3%R(2)%C = t2(6,5,4) +var3%R(2)%D = reshape([t2(1,2,3), t2(4,5,6), t2(11,12,13), t2(14,15,16)], [2,2]) + +! Assign + allocate +var4%Q%A = 145 +var4%Q%B = [991,992,993] +var4%Q%C = t2(996,995,994) +var4%Q%D = reshape([t2(199,299,399), t2(499,599,699), t2(1199,1299,1399), t2(1499,1599,1699)], [2,2]) + +var4%R(3)%A = 145 +var4%R(3)%B = [991,992,993] +var4%R(3)%C = t2(996,995,994) +var4%R(3)%D = reshape([t2(199,299,399), t2(499,599,699), t2(1199,1299,1399), t2(1499,1599,1699)], [2,2]) + +!$omp target map(to: var3%Q) map(tofrom: var4%Q) + call foo(var3%Q, var4%Q) +!$omp end target + +!$omp target map(to: var3%R(2)) map(tofrom: var4%R(3)) + call foo(var3%R(2), var4%R(3)) +!$omp end target + +if (var4%Q%A /= 45) stop 13 +if (any (var4%Q%B /= [1,2,3])) stop 14 +if (var4%Q%C%x /= 6) stop 15 +if (var4%Q%C%y /= 5) stop 15 +if (var4%Q%C%z /= 4) stop 15 +if (any (var4%Q%D(:,:)%x /= reshape([1, 4, 11, 14], [2,2]))) stop 16 +if (any (var4%Q%D(:,:)%y /= reshape([2, 5, 12, 15], [2,2]))) stop 16 +if (any (var4%Q%D(:,:)%z /= reshape([3, 6, 13, 16], [2,2]))) stop 16 + +if (var4%R(3)%A /= 45) stop 17 +if (any (var4%R(3)%B /= [1,2,3])) stop 18 +if (var4%R(3)%C%x /= 6) stop 19 +if (var4%R(3)%C%y /= 5) stop 19 +if (var4%R(3)%C%z /= 4) stop 19 +if (any (var4%R(3)%D(:,:)%x /= reshape([1, 4, 11, 14], [2,2]))) stop 20 +if (any (var4%R(3)%D(:,:)%y /= reshape([2, 5, 12, 15], [2,2]))) stop 20 +if (any (var4%R(3)%D(:,:)%z /= reshape([3, 6, 13, 16], [2,2]))) stop 20 + +contains + subroutine foo(x, y) + type(t) :: x, y + if (x%A /= 45) stop 1 + if (any (x%B /= [1,2,3])) stop 2 + if (x%C%x /= 6) stop 3 + if (x%C%y /= 5) stop 3 + if (x%C%z /= 4) stop 3 + if (any (x%D(:,:)%x /= reshape([1, 4, 11, 14], [2,2]))) stop 4 + if (any (x%D(:,:)%y /= reshape([2, 5, 12, 15], [2,2]))) stop 4 + if (any (x%D(:,:)%z /= reshape([3, 6, 13, 16], [2,2]))) stop 4 + + if (y%A /= 145) stop 5 + if (any (y%B /= [991,992,993])) stop 6 + if (y%C%x /= 996) stop 7 + if (y%C%y /= 995) stop 7 + if (y%C%z /= 994) stop 7 + if (any (y%D(:,:)%x /= reshape([199, 499, 1199, 1499], [2,2]))) stop 8 + if (any (y%D(:,:)%y /= reshape([299, 599, 1299, 1599], [2,2]))) stop 8 + if (any (y%D(:,:)%z /= reshape([399, 699, 1399, 1699], [2,2]))) stop 8 + + y%A = x%A + y%B(:) = x%B + y%C = x%C + y%D(:,:) = x%D(:,:) + end +end diff --git a/libgomp/testsuite/libgomp.fortran/map-alloc-comp-4.f90 b/libgomp/testsuite/libgomp.fortran/map-alloc-comp-4.f90 new file mode 100644 index 0000000..fb9859d --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/map-alloc-comp-4.f90 @@ -0,0 +1,124 @@ +type t2 + integer x, y, z +end type t2 +type t + integer, allocatable :: A + integer, allocatable :: B(:) + type(t2), allocatable :: C + type(t2), allocatable :: D(:,:) +end type t + +type t3 + type(t) :: Q + type(t) :: R(5) +end type + +type(t) :: var, var2 +type(t3) :: var3, var4 + +! -------------------------------------- +! Assign + allocate +var%A = 45 +var%B = [1,2,3] +var%C = t2(6,5,4) +var%D = reshape([t2(1,2,3), t2(4,5,6), t2(11,12,13), t2(14,15,16)], [2,2]) + +! Assign + allocate +var2%A = 145 +var2%B = [991,992,993] +var2%C = t2(996,995,994) +var2%D = reshape([t2(199,299,399), t2(499,599,699), t2(1199,1299,1399), t2(1499,1599,1699)], [2,2]) + + +!$omp target map(to: var%A, var%B, var%C, var%D) & +!$omp& map(tofrom: var2%A, var2%B, var2%C, var2%D) + call foo(var, var2) +!$omp end target + +if (var2%A /= 45) stop 9 +if (any (var2%B /= [1,2,3])) stop 10 +if (var2%C%x /= 6) stop 11 +if (var2%C%y /= 5) stop 11 +if (var2%C%z /= 4) stop 11 +if (any (var2%D(:,:)%x /= reshape([1, 4, 11, 14], [2,2]))) stop 12 +if (any (var2%D(:,:)%y /= reshape([2, 5, 12, 15], [2,2]))) stop 12 +if (any (var2%D(:,:)%z /= reshape([3, 6, 13, 16], [2,2]))) stop 12 + +! -------------------------------------- +! Assign + allocate +var3%Q%A = 45 +var3%Q%B = [1,2,3] +var3%Q%C = t2(6,5,4) +var3%Q%D = reshape([t2(1,2,3), t2(4,5,6), t2(11,12,13), t2(14,15,16)], [2,2]) + +var3%R(2)%A = 45 +var3%R(2)%B = [1,2,3] +var3%R(2)%C = t2(6,5,4) +var3%R(2)%D = reshape([t2(1,2,3), t2(4,5,6), t2(11,12,13), t2(14,15,16)], [2,2]) + +! Assign + allocate +var4%Q%A = 145 +var4%Q%B = [991,992,993] +var4%Q%C = t2(996,995,994) +var4%Q%D = reshape([t2(199,299,399), t2(499,599,699), t2(1199,1299,1399), t2(1499,1599,1699)], [2,2]) + +var4%R(3)%A = 145 +var4%R(3)%B = [991,992,993] +var4%R(3)%C = t2(996,995,994) +var4%R(3)%D = reshape([t2(199,299,399), t2(499,599,699), t2(1199,1299,1399), t2(1499,1599,1699)], [2,2]) + +!$omp target map(to: var3%Q%A, var3%Q%B, var3%Q%C, var3%Q%D) & +!$omp& map(tofrom: var4%Q%A, var4%Q%B, var4%Q%C, var4%Q%D) + call foo(var3%Q, var4%Q) +!$omp end target + +if (var4%Q%A /= 45) stop 13 +if (any (var4%Q%B /= [1,2,3])) stop 14 +if (var4%Q%C%x /= 6) stop 15 +if (var4%Q%C%y /= 5) stop 15 +if (var4%Q%C%z /= 4) stop 15 +if (any (var4%Q%D(:,:)%x /= reshape([1, 4, 11, 14], [2,2]))) stop 16 +if (any (var4%Q%D(:,:)%y /= reshape([2, 5, 12, 15], [2,2]))) stop 16 +if (any (var4%Q%D(:,:)%z /= reshape([3, 6, 13, 16], [2,2]))) stop 16 + +!$omp target map(to: var3%R(2)%A, var3%R(2)%B, var3%R(2)%C, var3%R(2)%D) & +!$omp& map(tofrom: var4%R(3)%A, var4%R(3)%B, var4%R(3)%C, var4%R(3)%D) + call foo(var3%R(2), var4%R(3)) +!$omp end target + +if (var4%R(3)%A /= 45) stop 17 +if (any (var4%R(3)%B /= [1,2,3])) stop 18 +if (var4%R(3)%C%x /= 6) stop 19 +if (var4%R(3)%C%y /= 5) stop 19 +if (var4%R(3)%C%z /= 4) stop 19 +if (any (var4%R(3)%D(:,:)%x /= reshape([1, 4, 11, 14], [2,2]))) stop 20 +if (any (var4%R(3)%D(:,:)%y /= reshape([2, 5, 12, 15], [2,2]))) stop 20 +if (any (var4%R(3)%D(:,:)%z /= reshape([3, 6, 13, 16], [2,2]))) stop 20 + +contains + subroutine foo(x, y) + type(t) :: x, y + if (x%A /= 45) stop 1 + if (any (x%B /= [1,2,3])) stop 2 + if (x%C%x /= 6) stop 3 + if (x%C%y /= 5) stop 3 + if (x%C%z /= 4) stop 3 + if (any (x%D(:,:)%x /= reshape([1, 4, 11, 14], [2,2]))) stop 4 + if (any (x%D(:,:)%y /= reshape([2, 5, 12, 15], [2,2]))) stop 4 + if (any (x%D(:,:)%z /= reshape([3, 6, 13, 16], [2,2]))) stop 4 + + if (y%A /= 145) stop 5 + if (any (y%B /= [991,992,993])) stop 6 + if (y%C%x /= 996) stop 7 + if (y%C%y /= 995) stop 7 + if (y%C%z /= 994) stop 7 + if (any (y%D(:,:)%x /= reshape([199, 499, 1199, 1499], [2,2]))) stop 8 + if (any (y%D(:,:)%y /= reshape([299, 599, 1299, 1599], [2,2]))) stop 8 + if (any (y%D(:,:)%z /= reshape([399, 699, 1399, 1699], [2,2]))) stop 8 + + y%A = x%A + y%B(:) = x%B + y%C = x%C + y%D(:,:) = x%D(:,:) + end +end diff --git a/libgomp/testsuite/libgomp.fortran/map-alloc-comp-5.f90 b/libgomp/testsuite/libgomp.fortran/map-alloc-comp-5.f90 new file mode 100644 index 0000000..b2e36b2 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/map-alloc-comp-5.f90 @@ -0,0 +1,53 @@ +implicit none +type t + integer, allocatable :: a, b(:) +end type t +type(t) :: x, y, z +integer :: i + +!$omp target + if (allocated(x%a)) stop 1 + if (allocated(x%b)) stop 2 +!$omp end target + +allocate(x%a, x%b(-4:6)) +x%b(:) = [(i, i=-4,6)] + +!$omp target + if (.not. allocated(x%a)) stop 3 + if (.not. allocated(x%b)) stop 4 + if (lbound(x%b,1) /= -4) stop 5 + if (ubound(x%b,1) /= 6) stop 6 + if (any (x%b /= [(i, i=-4,6)])) stop 7 +!$omp end target + + +! The following only works with arrays due to +! PR fortran/96668 + +!$omp target enter data map(to: y, z) + +!$omp target + if (allocated(y%b)) stop 8 + if (allocated(z%b)) stop 9 +!$omp end target + +allocate(y%b(5), z%b(3)) +y%b = 42 +z%b = 99 + +! (implicitly) 'tofrom' mapped +! Planned for OpenMP 6.0 (but common extension) +! OpenMP <= 5.0 unclear +!$omp target + if (.not.allocated(y%b)) stop 10 + if (any (y%b /= 42)) stop 11 +!$omp end target + +! always map: OpenMP 5.1 (clarified) +!$omp target map(always, tofrom: z) + if (.not.allocated(z%b)) stop 12 + if (any (z%b /= 99)) stop 13 +!$omp end target + +end diff --git a/libgomp/testsuite/libgomp.fortran/map-alloc-comp-6.f90 b/libgomp/testsuite/libgomp.fortran/map-alloc-comp-6.f90 new file mode 100644 index 0000000..48d4aea --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/map-alloc-comp-6.f90 @@ -0,0 +1,308 @@ +! NOTE: This code uses POINTER. +! While map(p, var%p) etc. maps the ptr/ptr comp p / var%p (incl. allocatable comps), +! map(var) does not map var%p. + +use iso_c_binding +implicit none +type t2 + integer, allocatable :: x, y, z +end type t2 +type t + integer, pointer :: A => null() + integer, pointer :: B(:) => null() + type(t2), pointer :: C => null() + type(t2), pointer :: D(:,:) => null() +end type t + +type t3 + type(t) :: Q + type(t) :: R(5) +end type + +type(t) :: var, var2 +type(t3) :: var3, var4 +integer(c_intptr_t) :: iptr + +! -------------------------------------- +! Assign + allocate +allocate (var%A, source=45) +allocate (var%B(3), source=[1,2,3]) +allocate (var%C) +var%C%x = 6; var%C%y = 5; var%C%z = 4 +allocate (var%D(2,2)) +var%D(1,1)%x = 1 +var%D(1,1)%y = 2 +var%D(1,1)%z = 3 +var%D(2,1)%x = 4 +var%D(2,1)%y = 5 +var%D(2,1)%z = 6 +var%D(1,2)%x = 11 +var%D(1,2)%y = 12 +var%D(1,2)%z = 13 +var%D(2,2)%x = 14 +var%D(2,2)%y = 15 +var%D(2,2)%z = 16 + +! Assign + allocate +allocate (var2%A, source=145) +allocate (var2%B, source=[991,992,993]) +allocate (var2%C) +var2%C%x = 996; var2%C%y = 995; var2%C%z = 994 +allocate (var2%D(2,2)) +var2%D(1,1)%x = 199 +var2%D(1,1)%y = 299 +var2%D(1,1)%z = 399 +var2%D(2,1)%x = 499 +var2%D(2,1)%y = 599 +var2%D(2,1)%z = 699 +var2%D(1,2)%x = 1199 +var2%D(1,2)%y = 1299 +var2%D(1,2)%z = 1399 +var2%D(2,2)%x = 1499 +var2%D(2,2)%y = 1599 +var2%D(2,2)%z = 1699 + +block + integer(c_intptr_t) :: loc_a, loc_b, loc_c, loc_d, loc2_a, loc2_b, loc2_c, loc2_d + loc_a = loc (var%a) + loc_b = loc (var%b) + loc_c = loc (var%d) + loc_d = loc (var%d) + loc2_a = loc (var2%a) + loc2_b = loc (var2%b) + loc2_c = loc (var2%c) + loc2_d = loc (var2%d) + ! var/var2 are mapped, but the pointer components aren't + !$omp target map(to: var) map(tofrom: var2) + if (loc_a /= loc (var%a)) stop 31 + if (loc_b /= loc (var%b)) stop 32 + if (loc_c /= loc (var%d)) stop 33 + if (loc_d /= loc (var%d)) stop 34 + if (loc2_a /= loc (var2%a)) stop 35 + if (loc2_b /= loc (var2%b)) stop 36 + if (loc2_c /= loc (var2%c)) stop 37 + if (loc2_d /= loc (var2%d)) stop 38 + !$omp end target + if (loc_a /= loc (var%a)) stop 41 + if (loc_b /= loc (var%b)) stop 42 + if (loc_c /= loc (var%d)) stop 43 + if (loc_d /= loc (var%d)) stop 44 + if (loc2_a /= loc (var2%a)) stop 45 + if (loc2_b /= loc (var2%b)) stop 46 + if (loc2_c /= loc (var2%c)) stop 47 + if (loc2_d /= loc (var2%d)) stop 48 +end block + +block + ! Map only (all) components, but this maps also the alloc comps + !$omp target map(to: var%a, var%b, var%c, var%d) map(tofrom: var2%a, var2%b, var2%c, var2%d) + call foo (var,var2) + !$omp end target +end block + +if (var2%A /= 45) stop 9 +if (any (var2%B /= [1,2,3])) stop 10 +if (var2%C%x /= 6) stop 11 +if (var2%C%y /= 5) stop 11 +if (var2%C%z /= 4) stop 11 +block + integer :: tmp_x(2,2), tmp_y(2,2), tmp_z(2,2), i, j + tmp_x = reshape([1, 4, 11, 14], [2,2]) + tmp_y = reshape([2, 5, 12, 15], [2,2]) + tmp_z = reshape([3, 6, 13, 16], [2,2]) + do j = 1, 2 + do i = 1, 2 + if (var2%D(i,j)%x /= tmp_x(i,j)) stop 12 + if (var2%D(i,j)%y /= tmp_y(i,j)) stop 12 + if (var2%D(i,j)%z /= tmp_z(i,j)) stop 12 + end do + end do +end block + +! Extra deallocates due to PR fortran/104697 +deallocate(var%C%x, var%C%y, var%C%z) +deallocate(var%D(1,1)%x, var%D(1,1)%y, var%D(1,1)%z) +deallocate(var%D(2,1)%x, var%D(2,1)%y, var%D(2,1)%z) +deallocate(var%D(1,2)%x, var%D(1,2)%y, var%D(1,2)%z) +deallocate(var%D(2,2)%x, var%D(2,2)%y, var%D(2,2)%z) +deallocate(var%A, var%B, var%C, var%D) + +deallocate(var2%C%x, var2%C%y, var2%C%z) +deallocate(var2%D(1,1)%x, var2%D(1,1)%y, var2%D(1,1)%z) +deallocate(var2%D(2,1)%x, var2%D(2,1)%y, var2%D(2,1)%z) +deallocate(var2%D(1,2)%x, var2%D(1,2)%y, var2%D(1,2)%z) +deallocate(var2%D(2,2)%x, var2%D(2,2)%y, var2%D(2,2)%z) +deallocate(var2%A, var2%B, var2%C, var2%D) + +! -------------------------------------- +! Assign + allocate +allocate (var3%Q%A, source=45) +allocate (var3%Q%B, source=[1,2,3]) +allocate (var3%Q%C, source=t2(6,5,4)) +allocate (var3%Q%D(2,2)) +var3%Q%D(1,1) = t2(1,2,3) +var3%Q%D(2,1) = t2(4,5,6) +var3%Q%D(1,2) = t2(11,12,13) +var3%Q%D(2,2) = t2(14,15,16) + +allocate (var3%R(2)%A, source=45) +allocate (var3%R(2)%B, source=[1,2,3]) +allocate (var3%R(2)%C, source=t2(6,5,4)) +allocate (var3%R(2)%D(2,2)) +var3%R(2)%D(1,1) = t2(1,2,3) +var3%R(2)%D(2,1) = t2(4,5,6) +var3%R(2)%D(1,2) = t2(11,12,13) +var3%R(2)%D(2,2) = t2(14,15,16) + +! Assign + allocate +allocate (var4%Q%A, source=145) +allocate (var4%Q%B, source=[991,992,993]) +allocate (var4%Q%C, source=t2(996,995,994)) +allocate (var4%Q%D(2,2)) +var4%Q%D(1,1) = t2(199,299,399) +var4%Q%D(2,1) = t2(499,599,699) +var4%Q%D(1,2) = t2(1199,1299,1399) +var4%Q%D(2,2) = t2(1499,1599,1699) + +allocate (var4%R(3)%A, source=145) +allocate (var4%R(3)%B, source=[991,992,993]) +allocate (var4%R(3)%C, source=t2(996,995,994)) +allocate (var4%R(3)%D(2,2)) +var4%R(3)%D(1,1) = t2(199,299,399) +var4%R(3)%D(2,1) = t2(499,599,699) +var4%R(3)%D(1,2) = t2(1199,1299,1399) +var4%R(3)%D(2,2) = t2(1499,1599,1699) + +!$omp target map(to: var3%Q%A, var3%Q%B, var3%Q%C, var3%Q%D) & +!$omp& map(tofrom: var4%Q%A, var4%Q%B, var4%Q%C, var4%Q%D) + call foo(var3%Q, var4%Q) +!$omp end target + +iptr = loc(var3%R(2)%A) + +!$omp target map(to: var3%R(2)%A, var3%R(2)%B, var3%R(2)%C, var3%R(2)%D) & +!$omp& map(tofrom: var4%R(3)%A, var4%R(3)%B, var4%R(3)%C, var4%R(3)%D) + call foo(var3%R(2), var4%R(3)) +!$omp end target + +if (var4%Q%A /= 45) stop 13 +if (any (var4%Q%B /= [1,2,3])) stop 14 +if (var4%Q%C%x /= 6) stop 15 +if (var4%Q%C%y /= 5) stop 15 +if (var4%Q%C%z /= 4) stop 15 +block + integer :: tmp_x(2,2), tmp_y(2,2), tmp_z(2,2), i, j + tmp_x = reshape([1, 4, 11, 14], [2,2]) + tmp_y = reshape([2, 5, 12, 15], [2,2]) + tmp_z = reshape([3, 6, 13, 16], [2,2]) + do j = 1, 2 + do i = 1, 2 + if (var4%Q%D(i,j)%x /= tmp_x(i,j)) stop 16 + if (var4%Q%D(i,j)%y /= tmp_y(i,j)) stop 16 + if (var4%Q%D(i,j)%z /= tmp_z(i,j)) stop 16 + end do + end do +end block + +! Cf. PR fortran/104696 +! { dg-output "valid mapping, OK" { xfail { offload_device_nonshared_as } } } +if (iptr /= loc(var3%R(2)%A)) then + print *, "invalid mapping, cf. PR fortran/104696" +else + +if (var4%R(3)%A /= 45) stop 17 +if (any (var4%R(3)%B /= [1,2,3])) stop 18 +if (var4%R(3)%C%x /= 6) stop 19 +if (var4%R(3)%C%y /= 5) stop 19 +if (var4%R(3)%C%z /= 4) stop 19 +block + integer :: tmp_x(2,2), tmp_y(2,2), tmp_z(2,2), i, j + tmp_x = reshape([1, 4, 11, 14], [2,2]) + tmp_y = reshape([2, 5, 12, 15], [2,2]) + tmp_z = reshape([3, 6, 13, 16], [2,2]) + do j = 1, 2 + do i = 1, 2 + if (var4%R(3)%D(i,j)%x /= tmp_x(i,j)) stop 20 + if (var4%R(3)%D(i,j)%y /= tmp_y(i,j)) stop 20 + if (var4%R(3)%D(i,j)%z /= tmp_z(i,j)) stop 20 + end do + end do +end block + +! Extra deallocates due to PR fortran/104697 +deallocate(var3%Q%C%x, var3%Q%D(1,1)%x, var3%Q%D(2,1)%x, var3%Q%D(1,2)%x, var3%Q%D(2,2)%x) +deallocate(var3%Q%C%y, var3%Q%D(1,1)%y, var3%Q%D(2,1)%y, var3%Q%D(1,2)%y, var3%Q%D(2,2)%y) +deallocate(var3%Q%C%z, var3%Q%D(1,1)%z, var3%Q%D(2,1)%z, var3%Q%D(1,2)%z, var3%Q%D(2,2)%z) +deallocate(var3%Q%A, var3%Q%B, var3%Q%C, var3%Q%D) + +deallocate(var4%Q%C%x, var4%Q%D(1,1)%x, var4%Q%D(2,1)%x, var4%Q%D(1,2)%x, var4%Q%D(2,2)%x) +deallocate(var4%Q%C%y, var4%Q%D(1,1)%y, var4%Q%D(2,1)%y, var4%Q%D(1,2)%y, var4%Q%D(2,2)%y) +deallocate(var4%Q%C%z, var4%Q%D(1,1)%z, var4%Q%D(2,1)%z, var4%Q%D(1,2)%z, var4%Q%D(2,2)%z) +deallocate(var4%Q%A, var4%Q%B, var4%Q%C, var4%Q%D) + +deallocate(var3%R(2)%C%x, var3%R(2)%D(1,1)%x, var3%R(2)%D(2,1)%x, var3%R(2)%D(1,2)%x, var3%R(2)%D(2,2)%x) +deallocate(var3%R(2)%C%y, var3%R(2)%D(1,1)%y, var3%R(2)%D(2,1)%y, var3%R(2)%D(1,2)%y, var3%R(2)%D(2,2)%y) +deallocate(var3%R(2)%C%z, var3%R(2)%D(1,1)%z, var3%R(2)%D(2,1)%z, var3%R(2)%D(1,2)%z, var3%R(2)%D(2,2)%z) +deallocate(var3%R(2)%A, var3%R(2)%B, var3%R(2)%C, var3%R(2)%D) + +deallocate(var4%R(3)%C%x, var4%R(3)%D(1,1)%x, var4%R(3)%D(2,1)%x, var4%R(3)%D(1,2)%x, var4%R(3)%D(2,2)%x) +deallocate(var4%R(3)%C%y, var4%R(3)%D(1,1)%y, var4%R(3)%D(2,1)%y, var4%R(3)%D(1,2)%y, var4%R(3)%D(2,2)%y) +deallocate(var4%R(3)%C%z, var4%R(3)%D(1,1)%z, var4%R(3)%D(2,1)%z, var4%R(3)%D(1,2)%z, var4%R(3)%D(2,2)%z) +deallocate(var4%R(3)%A, var4%R(3)%B, var4%R(3)%C, var4%R(3)%D) + + print *, "valid mapping, OK" +endif + +contains + subroutine foo(x, y) + type(t) :: x, y + intent(in) :: x + intent(inout) :: y + integer :: tmp_x(2,2), tmp_y(2,2), tmp_z(2,2), i, j + if (x%A /= 45) stop 1 + if (any (x%B /= [1,2,3])) stop 2 + if (x%C%x /= 6) stop 3 + if (x%C%y /= 5) stop 3 + if (x%C%z /= 4) stop 3 + + tmp_x = reshape([1, 4, 11, 14], [2,2]) + tmp_y = reshape([2, 5, 12, 15], [2,2]) + tmp_z = reshape([3, 6, 13, 16], [2,2]) + do j = 1, 2 + do i = 1, 2 + if (x%D(i,j)%x /= tmp_x(i,j)) stop 4 + if (x%D(i,j)%y /= tmp_y(i,j)) stop 4 + if (x%D(i,j)%z /= tmp_z(i,j)) stop 4 + end do + end do + + if (y%A /= 145) stop 5 + if (any (y%B /= [991,992,993])) stop 6 + if (y%C%x /= 996) stop 7 + if (y%C%y /= 995) stop 7 + if (y%C%z /= 994) stop 7 + tmp_x = reshape([199, 499, 1199, 1499], [2,2]) + tmp_y = reshape([299, 599, 1299, 1599], [2,2]) + tmp_z = reshape([399, 699, 1399, 1699], [2,2]) + do j = 1, 2 + do i = 1, 2 + if (y%D(i,j)%x /= tmp_x(i,j)) stop 8 + if (y%D(i,j)%y /= tmp_y(i,j)) stop 8 + if (y%D(i,j)%z /= tmp_z(i,j)) stop 8 + end do + end do + + y%A = x%A + y%B(:) = x%B + y%C%x = x%C%x + y%C%y = x%C%y + y%C%z = x%C%z + do j = 1, 2 + do i = 1, 2 + y%D(i,j)%x = x%D(i,j)%x + y%D(i,j)%y = x%D(i,j)%y + y%D(i,j)%z = x%D(i,j)%z + end do + end do + end +end diff --git a/libgomp/testsuite/libgomp.fortran/map-alloc-comp-7.f90 b/libgomp/testsuite/libgomp.fortran/map-alloc-comp-7.f90 new file mode 100644 index 0000000..1493c5f --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/map-alloc-comp-7.f90 @@ -0,0 +1,672 @@ +module m + implicit none (type, external) + type t + integer, allocatable :: arr(:,:) + integer :: var + integer, allocatable :: slr + end type t + +contains + + subroutine check_it (is_present, dummy_alloced, inner_alloc, & + scalar, array, a_scalar, a_array, & + l_scalar, l_array, la_scalar, la_array, & + opt_scalar, opt_array, a_opt_scalar, a_opt_array) + type(t), intent(inout) :: & + scalar, array(:,:), opt_scalar, opt_array(:,:), a_scalar, a_array(:,:), & + a_opt_scalar, a_opt_array(:,:), & + l_scalar, l_array(:,:), la_scalar, la_array(:,:) + optional :: opt_scalar, opt_array, a_opt_scalar, a_opt_array + allocatable :: a_scalar, a_array, a_opt_scalar, a_opt_array, la_scalar, la_array + logical, value :: is_present, dummy_alloced, inner_alloc + integer :: i, j, k, l + + ! CHECK VALUE + if (scalar%var /= 42) stop 1 + if (l_scalar%var /= 42) stop 1 + if (is_present) then + if (opt_scalar%var /= 42) stop 2 + end if + if (any (shape(array) /= [3,2])) stop 1 + if (any (shape(l_array) /= [3,2])) stop 1 + if (is_present) then + if (any (shape(opt_array) /= [3,2])) stop 1 + end if + do j = 1, 2 + do i = 1, 3 + if (array(i,j)%var /= i*97 + 100*41*j) stop 3 + if (l_array(i,j)%var /= i*97 + 100*41*j) stop 3 + if (is_present) then + if (opt_array(i,j)%var /= i*97 + 100*41*j) stop 4 + end if + end do + end do + + if (dummy_alloced) then + if (a_scalar%var /= 42) stop 1 + if (la_scalar%var /= 42) stop 1 + if (is_present) then + if (a_opt_scalar%var /= 42) stop 1 + end if + if (any (shape(a_array) /= [3,2])) stop 1 + if (any (shape(la_array) /= [3,2])) stop 1 + if (is_present) then + if (any (shape(a_opt_array) /= [3,2])) stop 1 + end if + do j = 1, 2 + do i = 1, 3 + if (a_array(i,j)%var /= i*97 + 100*41*j) stop 1 + if (la_array(i,j)%var /= i*97 + 100*41*j) stop 1 + if (is_present) then + if (a_opt_array(i,j)%var /= i*97 + 100*41*j) stop 1 + end if + end do + end do + else + if (allocated (a_scalar)) stop 1 + if (allocated (la_scalar)) stop 1 + if (allocated (a_array)) stop 1 + if (allocated (la_array)) stop 1 + if (is_present) then + if (allocated (a_opt_scalar)) stop 1 + if (allocated (a_opt_array)) stop 1 + end if + end if + + if (inner_alloc) then + if (scalar%slr /= 467) stop 5 + if (l_scalar%slr /= 467) stop 5 + if (a_scalar%slr /= 467) stop 6 + if (la_scalar%slr /= 467) stop 6 + if (is_present) then + if (opt_scalar%slr /= 467) stop 7 + if (a_opt_scalar%slr /= 467) stop 8 + end if + do j = 1, 2 + do i = 1, 3 + if (array(i,j)%slr /= (i*97 + 100*41*j) + 467) stop 9 + if (l_array(i,j)%slr /= (i*97 + 100*41*j) + 467) stop 9 + if (a_array(i,j)%slr /= (i*97 + 100*41*j) + 467) stop 10 + if (la_array(i,j)%slr /= (i*97 + 100*41*j) + 467) stop 10 + if (is_present) then + if (opt_array(i,j)%slr /= (i*97 + 100*41*j) + 467) stop 11 + if (a_opt_array(i,j)%slr /= (i*97 + 100*41*j) + 467) stop 12 + end if + end do + end do + + do l = 1, 5 + do k = 1, 4 + if (any (shape(scalar%arr) /= [4,5])) stop 1 + if (any (shape(l_scalar%arr) /= [4,5])) stop 1 + if (any (shape(a_scalar%arr) /= [4,5])) stop 1 + if (any (shape(la_scalar%arr) /= [4,5])) stop 1 + if (scalar%arr(k,l) /= (i*27 + 1000*11*j) + 467) stop 13 + if (l_scalar%arr(k,l) /= (i*27 + 1000*11*j) + 467) stop 13 + if (a_scalar%arr(k,l) /= (i*27 + 1000*11*j) + 467) stop 14 + if (la_scalar%arr(k,l) /= (i*27 + 1000*11*j) + 467) stop 14 + if (is_present) then + if (any (shape(opt_scalar%arr) /= [4,5])) stop 1 + if (any (shape(a_opt_scalar%arr) /= [4,5])) stop 1 + if (opt_scalar%arr(k,l) /= (i*27 + 1000*11*j) + 467) stop 15 + if (a_opt_scalar%arr(k,l) /= (i*27 + 1000*11*j) + 467) stop 16 + end if + end do + end do + do j = 1, 2 + do i = 1, 3 + if (any (shape(array(i,j)%arr) /= [i,j])) stop 1 + if (any (shape(l_array(i,j)%arr) /= [i,j])) stop 1 + if (any (shape(a_array(i,j)%arr) /= [i,j])) stop 1 + if (any (shape(la_array(i,j)%arr) /= [i,j])) stop 1 + if (is_present) then + if (any (shape(opt_array(i,j)%arr) /= [i,j])) stop 1 + if (any (shape(a_opt_array(i,j)%arr) /= [i,j])) stop 1 + endif + do l = 1, j + do k = 1, i + if (array(i,j)%arr(k,l) /= i*27 + 1000*11*j + 467 + 3*k +53*l) stop 17 + if (l_array(i,j)%arr(k,l) /= i*27 + 1000*11*j + 467 + 3*k +53*l) stop 17 + if (a_array(i,j)%arr(k,l) /= i*27 + 1000*11*j + 467 + 3*k +53*l) stop 18 + if (la_array(i,j)%arr(k,l) /= i*27 + 1000*11*j + 467 + 3*k +53*l) stop 18 + if (is_present) then + if (opt_array(i,j)%arr(k,l) /= i*27 + 1000*11*j + 467 + 3*k +53*l) stop 19 + if (a_opt_array(i,j)%arr(k,l) /= i*27 + 1000*11*j + 467 + 3*k +53*l) stop 20 + end if + end do + end do + end do + end do + else if (dummy_alloced) then + if (allocated (scalar%slr)) stop 1 + if (allocated (l_scalar%slr)) stop 1 + if (allocated (a_scalar%slr)) stop 1 + if (allocated (la_scalar%slr)) stop 1 + if (is_present) then + if (allocated (opt_scalar%slr)) stop 1 + if (allocated (a_opt_scalar%slr)) stop 1 + endif + if (allocated (scalar%arr)) stop 1 + if (allocated (l_scalar%arr)) stop 1 + if (allocated (a_scalar%arr)) stop 1 + if (allocated (la_scalar%arr)) stop 1 + if (is_present) then + if (allocated (opt_scalar%arr)) stop 1 + if (allocated (a_opt_scalar%arr)) stop 1 + endif + end if + + ! SET VALUE + scalar%var = 42 + 13 + l_scalar%var = 42 + 13 + if (is_present) then + opt_scalar%var = 42 + 13 + endif + do j = 1, 2 + do i = 1, 3 + array(i,j)%var = i*97 + 100*41*j + 13 + l_array(i,j)%var = i*97 + 100*41*j + 13 + if (is_present) then + opt_array(i,j)%var = i*97 + 100*41*j + 13 + end if + end do + end do + + if (dummy_alloced) then + a_scalar%var = 42 + 13 + la_scalar%var = 42 + 13 + if (is_present) then + a_opt_scalar%var = 42 + 13 + endif + do j = 1, 2 + do i = 1, 3 + a_array(i,j)%var = i*97 + 100*41*j + 13 + la_array(i,j)%var = i*97 + 100*41*j + 13 + if (is_present) then + a_opt_array(i,j)%var = i*97 + 100*41*j + 13 + endif + end do + end do + end if + + if (inner_alloc) then + scalar%slr = 467 + 13 + l_scalar%slr = 467 + 13 + a_scalar%slr = 467 + 13 + la_scalar%slr = 467 + 13 + if (is_present) then + opt_scalar%slr = 467 + 13 + a_opt_scalar%slr = 467 + 13 + end if + do j = 1, 2 + do i = 1, 3 + array(i,j)%slr = (i*97 + 100*41*j) + 467 + 13 + l_array(i,j)%slr = (i*97 + 100*41*j) + 467 + 13 + a_array(i,j)%slr = (i*97 + 100*41*j) + 467 + 13 + la_array(i,j)%slr = (i*97 + 100*41*j) + 467 + 13 + if (is_present) then + opt_array(i,j)%slr = (i*97 + 100*41*j) + 467 + 13 + a_opt_array(i,j)%slr = (i*97 + 100*41*j) + 467 + 13 + end if + end do + end do + + do l = 1, 5 + do k = 1, 4 + scalar%arr(k,l) = (i*27 + 1000*11*j) + 467 + 13 + l_scalar%arr(k,l) = (i*27 + 1000*11*j) + 467 + 13 + a_scalar%arr(k,l) = (i*27 + 1000*11*j) + 467 + 13 + la_scalar%arr(k,l) = (i*27 + 1000*11*j) + 467 + 13 + if (is_present) then + opt_scalar%arr(k,l) = (i*27 + 1000*11*j) + 467 + 13 + a_opt_scalar%arr(k,l) = (i*27 + 1000*11*j) + 467 + 13 + end if + end do + end do + do j = 1, 2 + do i = 1, 3 + do l = 1, j + do k = 1, i + array(i,j)%arr(k,l) = i*27 + 1000*11*j + 467 + 3*k +53*l + 13 + l_array(i,j)%arr(k,l) = i*27 + 1000*11*j + 467 + 3*k +53*l + 13 + a_array(i,j)%arr(k,l) = i*27 + 1000*11*j + 467 + 3*k +53*l + 13 + la_array(i,j)%arr(k,l) = i*27 + 1000*11*j + 467 + 3*k +53*l + 13 + if (is_present) then + opt_array(i,j)%arr(k,l) = i*27 + 1000*11*j + 467 + 3*k +53*l + 13 + a_opt_array(i,j)%arr(k,l) = i*27 + 1000*11*j + 467 + 3*k +53*l + 13 + end if + end do + end do + end do + end do + end if + + end subroutine + subroutine check_reset (is_present, dummy_alloced, inner_alloc, & + scalar, array, a_scalar, a_array, & + l_scalar, l_array, la_scalar, la_array, & + opt_scalar, opt_array, a_opt_scalar, a_opt_array) + type(t), intent(inout) :: & + scalar, array(:,:), opt_scalar, opt_array(:,:), a_scalar, a_array(:,:), & + a_opt_scalar, a_opt_array(:,:), & + l_scalar, l_array(:,:), la_scalar, la_array(:,:) + optional :: opt_scalar, opt_array, a_opt_scalar, a_opt_array + allocatable :: a_scalar, a_array, a_opt_scalar, a_opt_array, la_scalar, la_array + logical, value :: is_present, dummy_alloced, inner_alloc + integer :: i, j, k, l + + ! CHECK VALUE + if (scalar%var /= 42 + 13) stop 1 + if (l_scalar%var /= 42 + 13) stop 1 + if (is_present) then + if (opt_scalar%var /= 42 + 13) stop 2 + end if + if (any (shape(array) /= [3,2])) stop 1 + if (any (shape(l_array) /= [3,2])) stop 1 + if (is_present) then + if (any (shape(opt_array) /= [3,2])) stop 1 + end if + do j = 1, 2 + do i = 1, 3 + if (array(i,j)%var /= i*97 + 100*41*j + 13) stop 3 + if (l_array(i,j)%var /= i*97 + 100*41*j + 13) stop 3 + if (is_present) then + if (opt_array(i,j)%var /= i*97 + 100*41*j + 13) stop 4 + end if + end do + end do + + if (dummy_alloced) then + if (a_scalar%var /= 42 + 13) stop 1 + if (la_scalar%var /= 42 + 13) stop 1 + if (is_present) then + if (a_opt_scalar%var /= 42 + 13) stop 1 + end if + if (any (shape(a_array) /= [3,2])) stop 1 + if (any (shape(la_array) /= [3,2])) stop 1 + if (is_present) then + if (any (shape(a_opt_array) /= [3,2])) stop 1 + end if + do j = 1, 2 + do i = 1, 3 + if (a_array(i,j)%var /= i*97 + 100*41*j + 13) stop 1 + if (la_array(i,j)%var /= i*97 + 100*41*j + 13) stop 1 + if (is_present) then + if (a_opt_array(i,j)%var /= i*97 + 100*41*j + 13) stop 1 + end if + end do + end do + else + if (allocated (a_scalar)) stop 1 + if (allocated (la_scalar)) stop 1 + if (allocated (a_array)) stop 1 + if (allocated (la_array)) stop 1 + if (is_present) then + if (allocated (a_opt_scalar)) stop 1 + if (allocated (a_opt_array)) stop 1 + end if + end if + + if (inner_alloc) then + if (scalar%slr /= 467 + 13) stop 5 + if (l_scalar%slr /= 467 + 13) stop 5 + if (a_scalar%slr /= 467 + 13) stop 6 + if (la_scalar%slr /= 467 + 13) stop 6 + if (is_present) then + if (opt_scalar%slr /= 467 + 13) stop 7 + if (a_opt_scalar%slr /= 467 + 13) stop 8 + end if + do j = 1, 2 + do i = 1, 3 + if (array(i,j)%slr /= (i*97 + 100*41*j) + 467 + 13) stop 9 + if (l_array(i,j)%slr /= (i*97 + 100*41*j) + 467 + 13) stop 9 + if (a_array(i,j)%slr /= (i*97 + 100*41*j) + 467 + 13) stop 10 + if (la_array(i,j)%slr /= (i*97 + 100*41*j) + 467 + 13) stop 10 + if (is_present) then + if (opt_array(i,j)%slr /= (i*97 + 100*41*j) + 467 + 13) stop 11 + if (a_opt_array(i,j)%slr /= (i*97 + 100*41*j) + 467 + 13) stop 12 + end if + end do + end do + + do l = 1, 5 + do k = 1, 4 + if (any (shape(scalar%arr) /= [4,5])) stop 1 + if (any (shape(l_scalar%arr) /= [4,5])) stop 1 + if (any (shape(a_scalar%arr) /= [4,5])) stop 1 + if (any (shape(la_scalar%arr) /= [4,5])) stop 1 + if (scalar%arr(k,l) /= (i*27 + 1000*11*j) + 467 + 13) stop 13 + if (l_scalar%arr(k,l) /= (i*27 + 1000*11*j) + 467 + 13) stop 13 + if (a_scalar%arr(k,l) /= (i*27 + 1000*11*j) + 467 + 13) stop 14 + if (la_scalar%arr(k,l) /= (i*27 + 1000*11*j) + 467 + 13) stop 14 + if (is_present) then + if (any (shape(opt_scalar%arr) /= [4,5])) stop 1 + if (any (shape(a_opt_scalar%arr) /= [4,5])) stop 1 + if (opt_scalar%arr(k,l) /= (i*27 + 1000*11*j) + 467 + 13) stop 15 + if (a_opt_scalar%arr(k,l) /= (i*27 + 1000*11*j) + 467 + 13) stop 16 + end if + end do + end do + do j = 1, 2 + do i = 1, 3 + if (any (shape(array(i,j)%arr) /= [i,j])) stop 1 + if (any (shape(l_array(i,j)%arr) /= [i,j])) stop 1 + if (any (shape(a_array(i,j)%arr) /= [i,j])) stop 1 + if (any (shape(la_array(i,j)%arr) /= [i,j])) stop 1 + if (is_present) then + if (any (shape(opt_array(i,j)%arr) /= [i,j])) stop 1 + if (any (shape(a_opt_array(i,j)%arr) /= [i,j])) stop 1 + endif + do l = 1, j + do k = 1, i + if (array(i,j)%arr(k,l) /= i*27 + 1000*11*j + 467 + 3*k +53*l + 13) stop 17 + if (l_array(i,j)%arr(k,l) /= i*27 + 1000*11*j + 467 + 3*k +53*l + 13) stop 17 + if (a_array(i,j)%arr(k,l) /= i*27 + 1000*11*j + 467 + 3*k +53*l + 13) stop 18 + if (la_array(i,j)%arr(k,l) /= i*27 + 1000*11*j + 467 + 3*k +53*l + 13) stop 18 + if (is_present) then + if (opt_array(i,j)%arr(k,l) /= i*27 + 1000*11*j + 467 + 3*k +53*l + 13) stop 19 + if (a_opt_array(i,j)%arr(k,l) /= i*27 + 1000*11*j + 467 + 3*k +53*l + 13) stop 20 + end if + end do + end do + end do + end do + else if (dummy_alloced) then + if (allocated (scalar%slr)) stop 1 + if (allocated (l_scalar%slr)) stop 1 + if (allocated (a_scalar%slr)) stop 1 + if (allocated (la_scalar%slr)) stop 1 + if (is_present) then + if (allocated (opt_scalar%slr)) stop 1 + if (allocated (a_opt_scalar%slr)) stop 1 + endif + if (allocated (scalar%arr)) stop 1 + if (allocated (l_scalar%arr)) stop 1 + if (allocated (a_scalar%arr)) stop 1 + if (allocated (la_scalar%arr)) stop 1 + if (is_present) then + if (allocated (opt_scalar%arr)) stop 1 + if (allocated (a_opt_scalar%arr)) stop 1 + endif + end if + + ! (RE)SET VALUE + scalar%var = 42 + l_scalar%var = 42 + if (is_present) then + opt_scalar%var = 42 + endif + do j = 1, 2 + do i = 1, 3 + array(i,j)%var = i*97 + 100*41*j + l_array(i,j)%var = i*97 + 100*41*j + if (is_present) then + opt_array(i,j)%var = i*97 + 100*41*j + end if + end do + end do + + if (dummy_alloced) then + a_scalar%var = 42 + la_scalar%var = 42 + if (is_present) then + a_opt_scalar%var = 42 + endif + do j = 1, 2 + do i = 1, 3 + a_array(i,j)%var = i*97 + 100*41*j + la_array(i,j)%var = i*97 + 100*41*j + if (is_present) then + a_opt_array(i,j)%var = i*97 + 100*41*j + endif + end do + end do + end if + + if (inner_alloc) then + scalar%slr = 467 + l_scalar%slr = 467 + a_scalar%slr = 467 + la_scalar%slr = 467 + if (is_present) then + opt_scalar%slr = 467 + a_opt_scalar%slr = 467 + end if + do j = 1, 2 + do i = 1, 3 + array(i,j)%slr = (i*97 + 100*41*j) + 467 + l_array(i,j)%slr = (i*97 + 100*41*j) + 467 + a_array(i,j)%slr = (i*97 + 100*41*j) + 467 + la_array(i,j)%slr = (i*97 + 100*41*j) + 467 + if (is_present) then + opt_array(i,j)%slr = (i*97 + 100*41*j) + 467 + a_opt_array(i,j)%slr = (i*97 + 100*41*j) + 467 + end if + end do + end do + + do l = 1, 5 + do k = 1, 4 + scalar%arr(k,l) = (i*27 + 1000*11*j) + 467 + l_scalar%arr(k,l) = (i*27 + 1000*11*j) + 467 + a_scalar%arr(k,l) = (i*27 + 1000*11*j) + 467 + la_scalar%arr(k,l) = (i*27 + 1000*11*j) + 467 + if (is_present) then + opt_scalar%arr(k,l) = (i*27 + 1000*11*j) + 467 + a_opt_scalar%arr(k,l) = (i*27 + 1000*11*j) + 467 + end if + end do + end do + do j = 1, 2 + do i = 1, 3 + do l = 1, j + do k = 1, i + array(i,j)%arr(k,l) = i*27 + 1000*11*j + 467 + 3*k +53*l + l_array(i,j)%arr(k,l) = i*27 + 1000*11*j + 467 + 3*k +53*l + a_array(i,j)%arr(k,l) = i*27 + 1000*11*j + 467 + 3*k +53*l + la_array(i,j)%arr(k,l) = i*27 + 1000*11*j + 467 + 3*k +53*l + if (is_present) then + opt_array(i,j)%arr(k,l) = i*27 + 1000*11*j + 467 + 3*k +53*l + a_opt_array(i,j)%arr(k,l) = i*27 + 1000*11*j + 467 + 3*k +53*l + end if + end do + end do + end do + end do + end if + end subroutine + + subroutine test(scalar, array, a_scalar, a_array, opt_scalar, opt_array, & + a_opt_scalar, a_opt_array) + type(t) :: scalar, array(:,:), opt_scalar, opt_array(:,:), a_scalar, a_array(:,:) + type(t) :: a_opt_scalar, a_opt_array(:,:) + type(t) :: l_scalar, l_array(3,2), la_scalar, la_array(:,:) + allocatable :: a_scalar, a_array, a_opt_scalar, a_opt_array, la_scalar, la_array + optional :: opt_scalar, opt_array, a_opt_scalar, a_opt_array + + integer :: i, j, k, l + logical :: is_present, dummy_alloced, local_alloced, inner_alloc + is_present = present(opt_scalar) + dummy_alloced = allocated(a_scalar) + inner_alloc = allocated(scalar%slr) + + l_scalar%var = 42 + do j = 1, 2 + do i = 1, 3 + l_array(i,j)%var = i*97 + 100*41*j + end do + end do + + if (dummy_alloced) then + allocate(la_scalar, la_array(3,2)) + a_scalar%var = 42 + la_scalar%var = 42 + do j = 1, 2 + do i = 1, 3 + l_array(i,j)%var = i*97 + 100*41*j + la_array(i,j)%var = i*97 + 100*41*j + end do + end do + end if + + if (inner_alloc) then + l_scalar%slr = 467 + la_scalar%slr = 467 + do j = 1, 2 + do i = 1, 3 + l_array(i,j)%slr = (i*97 + 100*41*j) + 467 + la_array(i,j)%slr = (i*97 + 100*41*j) + 467 + end do + end do + + allocate(l_scalar%arr(4,5), la_scalar%arr(4,5)) + do l = 1, 5 + do k = 1, 4 + l_scalar%arr(k,l) = (i*27 + 1000*11*j) + 467 + la_scalar%arr(k,l) = (i*27 + 1000*11*j) + 467 + end do + end do + do j = 1, 2 + do i = 1, 3 + allocate(l_array(i,j)%arr(i,j), la_array(i,j)%arr(i,j)) + do l = 1, j + do k = 1, i + l_array(i,j)%arr(k,l) = i*27 + 1000*11*j + 467 + 3*k +53*l + la_array(i,j)%arr(k,l) = i*27 + 1000*11*j + 467 + 3*k +53*l + end do + end do + end do + end do + end if + + ! implicit mapping + !$omp target + if (is_present) then + call check_it (is_present, dummy_alloced, inner_alloc, & + scalar, array, a_scalar, a_array, & + l_scalar, l_array, la_scalar, la_array, & + opt_scalar, opt_array, a_opt_scalar, a_opt_array) + else + call check_it (is_present, dummy_alloced, inner_alloc, & + scalar, array, a_scalar, a_array, & + l_scalar, l_array, la_scalar, la_array) + end if + !$omp end target + + if (is_present) then + call check_reset (is_present, dummy_alloced, inner_alloc, & + scalar, array, a_scalar, a_array, & + l_scalar, l_array, la_scalar, la_array, & + opt_scalar, opt_array, a_opt_scalar, a_opt_array) + else + call check_reset (is_present, dummy_alloced, inner_alloc, & + scalar, array, a_scalar, a_array, & + l_scalar, l_array, la_scalar, la_array) + endif + + ! explicit mapping + !$omp target map(scalar, array, opt_scalar, opt_array, a_scalar, a_array) & + !$omp& map(a_opt_scalar, a_opt_array) & + !$omp& map(l_scalar, l_array, la_scalar, la_array) + if (is_present) then + call check_it (is_present, dummy_alloced, inner_alloc, & + scalar, array, a_scalar, a_array, & + l_scalar, l_array, la_scalar, la_array, & + opt_scalar, opt_array, a_opt_scalar, a_opt_array) + else + call check_it (is_present, dummy_alloced, inner_alloc, & + scalar, array, a_scalar, a_array, & + l_scalar, l_array, la_scalar, la_array) + endif + !$omp end target + + if (is_present) then + call check_reset (is_present, dummy_alloced, inner_alloc, & + scalar, array, a_scalar, a_array, & + l_scalar, l_array, la_scalar, la_array, & + opt_scalar, opt_array, a_opt_scalar, a_opt_array) + else + call check_reset (is_present, dummy_alloced, inner_alloc, & + scalar, array, a_scalar, a_array, & + l_scalar, l_array, la_scalar, la_array) + endif + end subroutine +end module + +program main + use m + implicit none (type, external) + type(t) :: scalar, array(3,2), opt_scalar, opt_array(3,2), a_scalar, a_array(:,:) + type(t) :: a_opt_scalar, a_opt_array(:,:) + allocatable :: a_scalar, a_array, a_opt_scalar, a_opt_array + integer :: i, j, k, l, n + + scalar%var = 42 + opt_scalar%var = 42 + do j = 1, 2 + do i = 1, 3 + array(i,j)%var = i*97 + 100*41*j + opt_array(i,j)%var = i*97 + 100*41*j + end do + end do + + ! unallocated + call test (scalar, array, a_scalar, a_array) + call test (scalar, array, a_scalar, a_array, opt_scalar, opt_array, a_opt_scalar, a_opt_array) + + ! allocated + allocate(a_scalar, a_opt_scalar, a_array(3,2), a_opt_array(3,2)) + a_scalar%var = 42 + a_opt_scalar%var = 42 + do j = 1, 2 + do i = 1, 3 + a_array(i,j)%var = i*97 + 100*41*j + a_opt_array(i,j)%var = i*97 + 100*41*j + end do + end do + + call test (scalar, array, a_scalar, a_array) + call test (scalar, array, a_scalar, a_array, opt_scalar, opt_array, a_opt_scalar, a_opt_array) + + ! comps allocated + scalar%slr = 467 + a_scalar%slr = 467 + opt_scalar%slr = 467 + a_opt_scalar%slr = 467 + do j = 1, 2 + do i = 1, 3 + array(i,j)%slr = (i*97 + 100*41*j) + 467 + a_array(i,j)%slr = (i*97 + 100*41*j) + 467 + opt_array(i,j)%slr = (i*97 + 100*41*j) + 467 + a_opt_array(i,j)%slr = (i*97 + 100*41*j) + 467 + end do + end do + + allocate(scalar%arr(4,5), a_scalar%arr(4,5), opt_scalar%arr(4,5), a_opt_scalar%arr(4,5)) + do l = 1, 5 + do k = 1, 4 + scalar%arr(k,l) = (i*27 + 1000*11*j) + 467 + a_scalar%arr(k,l) = (i*27 + 1000*11*j) + 467 + opt_scalar%arr(k,l) = (i*27 + 1000*11*j) + 467 + a_opt_scalar%arr(k,l) = (i*27 + 1000*11*j) + 467 + end do + end do + do j = 1, 2 + do i = 1, 3 + allocate(array(i,j)%arr(i,j), a_array(i,j)%arr(i,j), opt_array(i,j)%arr(i,j), a_opt_array(i,j)%arr(i,j)) + do l = 1, j + do k = 1, i + array(i,j)%arr(k,l) = i*27 + 1000*11*j + 467 + 3*k +53*l + a_array(i,j)%arr(k,l) = i*27 + 1000*11*j + 467 + 3*k +53*l + opt_array(i,j)%arr(k,l) = i*27 + 1000*11*j + 467 + 3*k +53*l + a_opt_array(i,j)%arr(k,l) = i*27 + 1000*11*j + 467 + 3*k +53*l + end do + end do + end do + end do + + call test (scalar, array, a_scalar, a_array) + call test (scalar, array, a_scalar, a_array, opt_scalar, opt_array, a_opt_scalar, a_opt_array) + + deallocate(a_scalar, a_opt_scalar, a_array, a_opt_array) +end diff --git a/libgomp/testsuite/libgomp.fortran/map-alloc-comp-8.f90 b/libgomp/testsuite/libgomp.fortran/map-alloc-comp-8.f90 new file mode 100644 index 0000000..f5a286e --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/map-alloc-comp-8.f90 @@ -0,0 +1,268 @@ +module m + implicit none (type, external) + type t + integer, allocatable :: A(:) + end type t + type t2 + type(t), allocatable :: vT + integer, allocatable :: x + end type t2 + +contains + + subroutine test_alloc() + type(t) :: var + type(t), allocatable :: var2 + + allocate(var2) + allocate(var%A(4), var2%A(5)) + + !$omp target enter data map(alloc: var, var2) + !$omp target + if (.not. allocated(Var2)) stop 1 + if (.not. allocated(Var%A)) stop 2 + if (.not. allocated(Var2%A)) stop 3 + if (lbound(var%A, 1) /= 1 .or. ubound(var%A, 1) /= 4) stop 4 + if (lbound(var2%A, 1) /= 1 .or. ubound(var2%A, 1) /= 5) stop 5 + var%A = [1,2,3,4] + var2%A = [11,22,33,44,55] + !$omp end target + !$omp target exit data map(from: var, var2) + + if (.not. allocated(Var2)) error stop + if (.not. allocated(Var%A)) error stop + if (.not. allocated(Var2%A)) error stop + if (lbound(var%A, 1) /= 1 .or. ubound(var%A, 1) /= 4) error stop + if (lbound(var2%A, 1) /= 1 .or. ubound(var2%A, 1) /= 5) error stop + if (any(var%A /= [1,2,3,4])) error stop + if (any(var2%A /= [11,22,33,44,55])) error stop + end subroutine test_alloc + + subroutine test2_alloc() + type(t2) :: var + type(t2), allocatable :: var2 + + allocate(var2) + allocate(var%x, var2%x) + + !$omp target enter data map(alloc: var, var2) + !$omp target + if (.not. allocated(Var2)) stop 6 + if (.not. allocated(Var%x)) stop 7 + if (.not. allocated(Var2%x)) stop 8 + var%x = 42 + var2%x = 43 + !$omp end target + !$omp target exit data map(from: var, var2) + + if (.not. allocated(Var2)) error stop + if (.not. allocated(Var%x)) error stop + if (.not. allocated(Var2%x)) error stop + if (var%x /= 42) error stop + if (var2%x /= 43) error stop + + allocate(var%vt, var2%vt) + allocate(var%vt%A(-1:3), var2%vt%A(0:4)) + + !$omp target enter data map(alloc: var, var2) + !$omp target + if (.not. allocated(Var2)) stop 11 + if (.not. allocated(Var%x)) stop 12 + if (.not. allocated(Var2%x)) stop 13 + if (.not. allocated(Var%vt)) stop 14 + if (.not. allocated(Var2%vt)) stop 15 + if (.not. allocated(Var%vt%a)) stop 16 + if (.not. allocated(Var2%vt%a)) stop 17 + var%x = 42 + var2%x = 43 + if (lbound(var%vt%A, 1) /= -1 .or. ubound(var%vt%A, 1) /= 3) stop 4 + if (lbound(var2%vt%A, 1) /= 0 .or. ubound(var2%vt%A, 1) /= 4) stop 5 + var%vt%A = [1,2,3,4,5] + var2%vt%A = [11,22,33,44,55] + !$omp end target + !$omp target exit data map(from: var, var2) + + if (.not. allocated(Var2)) error stop + if (.not. allocated(Var%x)) error stop + if (.not. allocated(Var2%x)) error stop + if (.not. allocated(Var%vt)) error stop + if (.not. allocated(Var2%vt)) error stop + if (.not. allocated(Var%vt%a)) error stop + if (.not. allocated(Var2%vt%a)) error stop + if (var%x /= 42) error stop + if (var2%x /= 43) error stop + if (lbound(var%vt%A, 1) /= -1 .or. ubound(var%vt%A, 1) /= 3) error stop + if (lbound(var2%vt%A, 1) /= 0 .or. ubound(var2%vt%A, 1) /= 4) error stop + if (any(var%vt%A /= [1,2,3,4,5])) error stop + if (any(var2%vt%A /= [11,22,33,44,55])) error stop + end subroutine test2_alloc + + + subroutine test_alloc_target() + type(t) :: var + type(t), allocatable :: var2 + + allocate(var2) + allocate(var%A(4), var2%A(5)) + + !$omp target map(alloc: var, var2) + if (.not. allocated(Var2)) stop 1 + if (.not. allocated(Var%A)) stop 2 + if (.not. allocated(Var2%A)) stop 3 + if (lbound(var%A, 1) /= 1 .or. ubound(var%A, 1) /= 4) stop 4 + if (lbound(var2%A, 1) /= 1 .or. ubound(var2%A, 1) /= 5) stop 5 + var%A = [1,2,3,4] + var2%A = [11,22,33,44,55] + !$omp end target + + if (.not. allocated(Var2)) error stop + if (.not. allocated(Var%A)) error stop + if (.not. allocated(Var2%A)) error stop + if (lbound(var%A, 1) /= 1 .or. ubound(var%A, 1) /= 4) error stop + if (lbound(var2%A, 1) /= 1 .or. ubound(var2%A, 1) /= 5) error stop + end subroutine test_alloc_target + + subroutine test2_alloc_target() + type(t2) :: var + type(t2), allocatable :: var2 + + allocate(var2) + allocate(var%x, var2%x) + + !$omp target map(alloc: var, var2) + if (.not. allocated(Var2)) stop 6 + if (.not. allocated(Var%x)) stop 7 + if (.not. allocated(Var2%x)) stop 8 + var%x = 42 + var2%x = 43 + !$omp end target + + if (.not. allocated(Var2)) error stop + if (.not. allocated(Var%x)) error stop + if (.not. allocated(Var2%x)) error stop + + allocate(var%vt, var2%vt) + allocate(var%vt%A(-1:3), var2%vt%A(0:4)) + + !$omp target map(alloc: var, var2) + if (.not. allocated(Var2)) stop 11 + if (.not. allocated(Var%x)) stop 12 + if (.not. allocated(Var2%x)) stop 13 + if (.not. allocated(Var%vt)) stop 14 + if (.not. allocated(Var2%vt)) stop 15 + if (.not. allocated(Var%vt%a)) stop 16 + if (.not. allocated(Var2%vt%a)) stop 17 + var%x = 42 + var2%x = 43 + if (lbound(var%vt%A, 1) /= -1 .or. ubound(var%vt%A, 1) /= 3) stop 4 + if (lbound(var2%vt%A, 1) /= 0 .or. ubound(var2%vt%A, 1) /= 4) stop 5 + var%vt%A = [1,2,3,4,5] + var2%vt%A = [11,22,33,44,55] + !$omp end target + + if (.not. allocated(Var2)) error stop + if (.not. allocated(Var%x)) error stop + if (.not. allocated(Var2%x)) error stop + if (.not. allocated(Var%vt)) error stop + if (.not. allocated(Var2%vt)) error stop + if (.not. allocated(Var%vt%a)) error stop + if (.not. allocated(Var2%vt%a)) error stop + if (lbound(var%vt%A, 1) /= -1 .or. ubound(var%vt%A, 1) /= 3) error stop + if (lbound(var2%vt%A, 1) /= 0 .or. ubound(var2%vt%A, 1) /= 4) error stop + end subroutine test2_alloc_target + + + + subroutine test_from() + type(t) :: var + type(t), allocatable :: var2 + + allocate(var2) + allocate(var%A(4), var2%A(5)) + + !$omp target map(from: var, var2) + if (.not. allocated(Var2)) stop 1 + if (.not. allocated(Var%A)) stop 2 + if (.not. allocated(Var2%A)) stop 3 + if (lbound(var%A, 1) /= 1 .or. ubound(var%A, 1) /= 4) stop 4 + if (lbound(var2%A, 1) /= 1 .or. ubound(var2%A, 1) /= 5) stop 5 + var%A = [1,2,3,4] + var2%A = [11,22,33,44,55] + !$omp end target + + if (.not. allocated(Var2)) error stop + if (.not. allocated(Var%A)) error stop + if (.not. allocated(Var2%A)) error stop + if (lbound(var%A, 1) /= 1 .or. ubound(var%A, 1) /= 4) error stop + if (lbound(var2%A, 1) /= 1 .or. ubound(var2%A, 1) /= 5) error stop + if (any(var%A /= [1,2,3,4])) error stop + if (any(var2%A /= [11,22,33,44,55])) error stop + end subroutine test_from + + subroutine test2_from() + type(t2) :: var + type(t2), allocatable :: var2 + + allocate(var2) + allocate(var%x, var2%x) + + !$omp target map(from: var, var2) + if (.not. allocated(Var2)) stop 6 + if (.not. allocated(Var%x)) stop 7 + if (.not. allocated(Var2%x)) stop 8 + var%x = 42 + var2%x = 43 + !$omp end target + + if (.not. allocated(Var2)) error stop + if (.not. allocated(Var%x)) error stop + if (.not. allocated(Var2%x)) error stop + if (var%x /= 42) error stop + if (var2%x /= 43) error stop + + allocate(var%vt, var2%vt) + allocate(var%vt%A(-1:3), var2%vt%A(0:4)) + + !$omp target map(from: var, var2) + if (.not. allocated(Var2)) stop 11 + if (.not. allocated(Var%x)) stop 12 + if (.not. allocated(Var2%x)) stop 13 + if (.not. allocated(Var%vt)) stop 14 + if (.not. allocated(Var2%vt)) stop 15 + if (.not. allocated(Var%vt%a)) stop 16 + if (.not. allocated(Var2%vt%a)) stop 17 + var%x = 42 + var2%x = 43 + if (lbound(var%vt%A, 1) /= -1 .or. ubound(var%vt%A, 1) /= 3) stop 4 + if (lbound(var2%vt%A, 1) /= 0 .or. ubound(var2%vt%A, 1) /= 4) stop 5 + var%vt%A = [1,2,3,4,5] + var2%vt%A = [11,22,33,44,55] + !$omp end target + + if (.not. allocated(Var2)) error stop + if (.not. allocated(Var%x)) error stop + if (.not. allocated(Var2%x)) error stop + if (.not. allocated(Var%vt)) error stop + if (.not. allocated(Var2%vt)) error stop + if (.not. allocated(Var%vt%a)) error stop + if (.not. allocated(Var2%vt%a)) error stop + if (var%x /= 42) error stop + if (var2%x /= 43) error stop + if (lbound(var%vt%A, 1) /= -1 .or. ubound(var%vt%A, 1) /= 3) error stop + if (lbound(var2%vt%A, 1) /= 0 .or. ubound(var2%vt%A, 1) /= 4) error stop + if (any(var%vt%A /= [1,2,3,4,5])) error stop + if (any(var2%vt%A /= [11,22,33,44,55])) error stop + end subroutine test2_from + +end module m + +use m + implicit none (type, external) + call test_alloc + call test2_alloc + call test_alloc_target + call test2_alloc_target + + call test_from + call test2_from +end diff --git a/libgomp/testsuite/libgomp.fortran/map-alloc-comp-9.f90 b/libgomp/testsuite/libgomp.fortran/map-alloc-comp-9.f90 new file mode 100644 index 0000000..3cec392 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/map-alloc-comp-9.f90 @@ -0,0 +1,559 @@ +! Ensure that polymorphic mapping is diagnosed as undefined behavior +! Ensure that static access to polymorphic variables works + +subroutine test(case) +implicit none(type, external) +type t + integer :: x(4) +end type t + +type ta + integer, allocatable :: x(:) +end type ta + +type t2 + class(t), allocatable :: x + class(t), allocatable :: x2(:) +end type t2 + +type t3 + type(t2) :: y + type(t2) :: y2(2) +end type t3 + +type t4 + type(t3), allocatable :: y + type(t3), allocatable :: y2(:) +end type t4 + +integer, value :: case + +logical :: is_shared_mem + +! Mangle stack addresses +integer, volatile :: case_var(100*case) + +type(t), allocatable :: var1 +type(ta), allocatable :: var1a +class(t), allocatable :: var2 +type(t2), allocatable :: var3 +type(t4), allocatable :: var4 + +case_var(100) = 0 +!print *, 'case', case + +var1 = t([1,2,3,4]) +var1a = ta([-1,-2,-3,-4,-5]) + +var2 = t([11,22,33,44]) + +allocate(t2 :: var3) +allocate(t :: var3%x) +allocate(t :: var3%x2(2)) +var3%x%x = [111,222,333,444] +var3%x2(1)%x = 2*[111,222,333,444] +var3%x2(2)%x = 3*[111,222,333,444] + +allocate(t4 :: var4) +allocate(t3 :: var4%y) +allocate(t3 :: var4%y2(2)) +allocate(t :: var4%y%y%x) +allocate(t :: var4%y%y%x2(2)) +allocate(t :: var4%y2(1)%y%x) +allocate(t :: var4%y2(1)%y%x2(2)) +allocate(t :: var4%y2(2)%y%x) +allocate(t :: var4%y2(2)%y%x2(2)) +var4%y%y%x%x = -1 * [1111,2222,3333,4444] +var4%y%y%x2(1)%x = -2 * [1111,2222,3333,4444] +var4%y%y%x2(2)%x = -3 * [1111,2222,3333,4444] +var4%y2(1)%y%x%x = -4 * [1111,2222,3333,4444] +var4%y2(1)%y%x2(1)%x = -5 * [1111,2222,3333,4444] +var4%y2(1)%y%x2(2)%x = -6 * [1111,2222,3333,4444] +var4%y2(2)%y%x%x = -7 * [1111,2222,3333,4444] +var4%y2(2)%y%x2(1)%x = -8 * [1111,2222,3333,4444] +var4%y2(2)%y%x2(2)%x = -9 * [1111,2222,3333,4444] + +is_shared_mem = .false. +!$omp target map(to: is_shared_mem) + is_shared_mem = .true. +!$omp end target + +if (case == 1) then + ! implicit mapping + !$omp target + if (any (var1%x /= [1,2,3,4])) stop 1 + var1%x = 2 * var1%x + !$omp end target + + !$omp target + if (any (var1a%x /= [-1,-2,-3,-4])) stop 2 + var1a%x = 3 * var1a%x + !$omp end target + + !$omp target ! { dg-warning "Mapping of polymorphic list item 'var2' is unspecified behavior \\\[-Wopenmp\\\]" } + if (any (var2%x /= [11,22,33,44])) stop 3 + var2%x = 4 * var2%x + !$omp end target + + !$omp target ! { dg-warning "Mapping of polymorphic list item 'var3->x' is unspecified behavior \\\[-Wopenmp\\\]" } + if (any (var3%x%x /= [111,222,333,444])) stop 4 + var3%x%x = 5 * var3%x%x + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + if (any (var3%x2(1)%x /= 2*[111,222,333,444])) stop 4 + if (any (var3%x2(2)%x /= 3*[111,222,333,444])) stop 4 + var3%x2(1)%x = 5 * var3%x2(1)%x + var3%x2(2)%x = 5 * var3%x2(2)%x + end if + !$omp end target + + !$omp target ! { dg-warning "Mapping of polymorphic list item 'var4\.\[0-9\]+->y->y\.x' is unspecified behavior \\\[-Wopenmp\\\]" } + if (any (var4%y%y%x%x /= -1 * [1111,2222,3333,4444])) stop 5 + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + if (any (var4%y%y%x2(1)%x /= -2 * [1111,2222,3333,4444])) stop 5 + if (any (var4%y%y%x2(2)%x /= -3 * [1111,2222,3333,4444])) stop 5 + endif + if (any (var4%y2(1)%y%x%x /= -4 * [1111,2222,3333,4444])) stop 5 + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + if (any (var4%y2(1)%y%x2(1)%x /= -5 * [1111,2222,3333,4444])) stop 5 + if (any (var4%y2(1)%y%x2(2)%x /= -6 * [1111,2222,3333,4444])) stop 5 + endif + if (any (var4%y2(2)%y%x%x /= -7 * [1111,2222,3333,4444])) stop 5 + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + if (any (var4%y2(2)%y%x2(1)%x /= -8 * [1111,2222,3333,4444])) stop 5 + if (any (var4%y2(2)%y%x2(2)%x /= -9 * [1111,2222,3333,4444])) stop 5 + end if + var4%y%y%x%x = 6 * var4%y%y%x%x + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + var4%y%y%x2(1)%x = 6 * var4%y%y%x2(1)%x + var4%y%y%x2(2)%x = 6 * var4%y%y%x2(2)%x + endif + var4%y2(1)%y%x%x = 6 * var4%y2(1)%y%x%x + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + var4%y2(1)%y%x2(1)%x = 6 * var4%y2(1)%y%x2(1)%x + var4%y2(1)%y%x2(2)%x = 6 * var4%y2(1)%y%x2(2)%x + endif + var4%y2(2)%y%x%x = 6 * var4%y2(2)%y%x%x + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + var4%y2(2)%y%x2(1)%x = 6 * var4%y2(2)%y%x2(1)%x + var4%y2(2)%y%x2(2)%x = 6 * var4%y2(2)%y%x2(2)%x + endif + !$omp end target + +else if (case == 2) then + ! Use target with defaultmap(TO) + + !$omp target defaultmap(to : all) + if (any (var1%x /= [1,2,3,4])) stop 1 + var1%x = 2 * var1%x + !$omp end target + + !$omp target defaultmap(to : all) + if (any (var1a%x /= [-1,-2,-3,-4])) stop 2 + var1a%x = 3 * var1a%x + !$omp end target + + !$omp target defaultmap(to : all) ! { dg-warning "Mapping of polymorphic list item 'var2' is unspecified behavior \\\[-Wopenmp\\\]" } + if (any (var2%x /= [11,22,33,44])) stop 3 + var2%x = 4 * var2%x + !$omp end target + + !$omp target defaultmap(to : all) ! { dg-warning "Mapping of polymorphic list item 'var3->x' is unspecified behavior \\\[-Wopenmp\\\]" } + if (any (var3%x%x /= [111,222,333,444])) stop 4 + var3%x%x = 5 * var3%x%x + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + if (any (var3%x2(1)%x /= 2*[111,222,333,444])) stop 4 + if (any (var3%x2(2)%x /= 3*[111,222,333,444])) stop 4 + var3%x2(1)%x = 5 * var3%x2(1)%x + var3%x2(2)%x = 5 * var3%x2(2)%x + endif + !$omp end target + + !$omp target defaultmap(to : all) firstprivate(is_shared_mem) ! { dg-warning "Mapping of polymorphic list item 'var4\.\[0-9\]+->y->y\.x' is unspecified behavior \\\[-Wopenmp\\\]" } + if (any (var4%y%y%x%x /= -1 * [1111,2222,3333,4444])) stop 5 + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + if (any (var4%y%y%x2(1)%x /= -2 * [1111,2222,3333,4444])) stop 5 + if (any (var4%y%y%x2(2)%x /= -3 * [1111,2222,3333,4444])) stop 5 + endif + if (any (var4%y2(1)%y%x%x /= -4 * [1111,2222,3333,4444])) stop 5 + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + if (any (var4%y2(1)%y%x2(1)%x /= -5 * [1111,2222,3333,4444])) stop 5 + if (any (var4%y2(1)%y%x2(2)%x /= -6 * [1111,2222,3333,4444])) stop 5 + endif + if (any (var4%y2(2)%y%x%x /= -7 * [1111,2222,3333,4444])) stop 5 + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + if (any (var4%y2(2)%y%x2(1)%x /= -8 * [1111,2222,3333,4444])) stop 5 + if (any (var4%y2(2)%y%x2(2)%x /= -9 * [1111,2222,3333,4444])) stop 5 + endif + var4%y%y%x%x = 6 * var4%y%y%x%x + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + var4%y%y%x2(1)%x = 6 * var4%y%y%x2(1)%x + var4%y%y%x2(2)%x = 6 * var4%y%y%x2(2)%x + endif + var4%y2(1)%y%x%x = 6 * var4%y2(1)%y%x%x + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + var4%y2(1)%y%x2(1)%x = 6 * var4%y2(1)%y%x2(1)%x + var4%y2(1)%y%x2(2)%x = 6 * var4%y2(1)%y%x2(2)%x + endif + var4%y2(2)%y%x%x = 6 * var4%y2(2)%y%x%x + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + var4%y2(2)%y%x2(1)%x = 6 * var4%y2(2)%y%x2(1)%x + var4%y2(2)%y%x2(2)%x = 6 * var4%y2(2)%y%x2(2)%x + endif + !$omp end target + +else if (case == 3) then + ! Use target with map clause + + !$omp target map(tofrom: var1) + if (any (var1%x /= [1,2,3,4])) stop 1 + var1%x = 2 * var1%x + !$omp end target + + !$omp target map(tofrom: var1a) + if (any (var1a%x /= [-1,-2,-3,-4])) stop 2 + var1a%x = 3 * var1a%x + !$omp end target + + !$omp target map(tofrom: var2) ! { dg-warning "28: Mapping of polymorphic list item 'var2' is unspecified behavior \\\[-Wopenmp\\\]" } + if (any (var2%x /= [11,22,33,44])) stop 3 + var2%x = 4 * var2%x + !$omp end target + + !$omp target map(tofrom: var3) ! { dg-warning "28: Mapping of polymorphic list item 'var3->x' is unspecified behavior \\\[-Wopenmp\\\]" } + if (any (var3%x%x /= [111,222,333,444])) stop 4 + var3%x%x = 5 * var3%x%x + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + if (any (var3%x2(1)%x /= 2*[111,222,333,444])) stop 4 + if (any (var3%x2(2)%x /= 3*[111,222,333,444])) stop 4 + var3%x2(1)%x = 5 * var3%x2(1)%x + var3%x2(2)%x = 5 * var3%x2(2)%x + endif + !$omp end target + + !$omp target map(tofrom: var4) ! { dg-warning "28: Mapping of polymorphic list item 'var4\.\[0-9\]+->y->y\.x' is unspecified behavior \\\[-Wopenmp\\\]" } + if (any (var4%y%y%x%x /= -1 * [1111,2222,3333,4444])) stop 5 + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + if (any (var4%y%y%x2(1)%x /= -2 * [1111,2222,3333,4444])) stop 5 + if (any (var4%y%y%x2(2)%x /= -3 * [1111,2222,3333,4444])) stop 5 + end if + if (any (var4%y2(1)%y%x%x /= -4 * [1111,2222,3333,4444])) stop 5 + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + if (any (var4%y2(1)%y%x2(1)%x /= -5 * [1111,2222,3333,4444])) stop 5 + if (any (var4%y2(1)%y%x2(2)%x /= -6 * [1111,2222,3333,4444])) stop 5 + endif + if (any (var4%y2(2)%y%x%x /= -7 * [1111,2222,3333,4444])) stop 5 + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + if (any (var4%y2(2)%y%x2(1)%x /= -8 * [1111,2222,3333,4444])) stop 5 + if (any (var4%y2(2)%y%x2(2)%x /= -9 * [1111,2222,3333,4444])) stop 5 + endif + var4%y%y%x%x = 6 * var4%y%y%x%x + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + var4%y%y%x2(1)%x = 6 * var4%y%y%x2(1)%x + var4%y%y%x2(2)%x = 6 * var4%y%y%x2(2)%x + endif + var4%y2(1)%y%x%x = 6 * var4%y2(1)%y%x%x + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + var4%y2(1)%y%x2(1)%x = 6 * var4%y2(1)%y%x2(1)%x + var4%y2(1)%y%x2(2)%x = 6 * var4%y2(1)%y%x2(2)%x + endif + var4%y2(2)%y%x%x = 6 * var4%y2(2)%y%x%x + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + var4%y2(2)%y%x2(1)%x = 6 * var4%y2(2)%y%x2(1)%x + var4%y2(2)%y%x2(2)%x = 6 * var4%y2(2)%y%x2(2)%x + endif + !$omp end target + +else if (case == 4) then + ! Use target with map clause -- NOTE: This uses TO not TOFROM + + !$omp target map(to: var1) + if (any (var1%x /= [1,2,3,4])) stop 1 + var1%x = 2 * var1%x + !$omp end target + + !$omp target map(to: var1a) + if (any (var1a%x /= [-1,-2,-3,-4])) stop 2 + var1a%x = 3 * var1a%x + !$omp end target + + !$omp target map(to: var2) ! { dg-warning "24: Mapping of polymorphic list item 'var2' is unspecified behavior \\\[-Wopenmp\\\]" } + if (any (var2%x /= [11,22,33,44])) stop 3 + var2%x = 4 * var2%x + !$omp end target + + !$omp target map(to: var3) ! { dg-warning "24: Mapping of polymorphic list item 'var3->x' is unspecified behavior \\\[-Wopenmp\\\]" } + if (any (var3%x%x /= [111,222,333,444])) stop 4 + var3%x%x = 5 * var3%x%x + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + if (any (var3%x2(1)%x /= 2*[111,222,333,444])) stop 4 + if (any (var3%x2(2)%x /= 3*[111,222,333,444])) stop 4 + var3%x2(1)%x = 5 * var3%x2(1)%x + var3%x2(2)%x = 5 * var3%x2(2)%x + endif + !$omp end target + + !$omp target map(to: var4) ! { dg-warning "24: Mapping of polymorphic list item 'var4\.\[0-9\]+->y->y\.x' is unspecified behavior \\\[-Wopenmp\\\]" } + if (any (var4%y%y%x%x /= -1 * [1111,2222,3333,4444])) stop 5 + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + if (any (var4%y%y%x2(1)%x /= -2 * [1111,2222,3333,4444])) stop 5 + if (any (var4%y%y%x2(2)%x /= -3 * [1111,2222,3333,4444])) stop 5 + endif + if (any (var4%y2(1)%y%x%x /= -4 * [1111,2222,3333,4444])) stop 5 + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + if (any (var4%y2(1)%y%x2(1)%x /= -5 * [1111,2222,3333,4444])) stop 5 + if (any (var4%y2(1)%y%x2(2)%x /= -6 * [1111,2222,3333,4444])) stop 5 + endif + if (any (var4%y2(2)%y%x%x /= -7 * [1111,2222,3333,4444])) stop 5 + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + if (any (var4%y2(2)%y%x2(1)%x /= -8 * [1111,2222,3333,4444])) stop 5 + if (any (var4%y2(2)%y%x2(2)%x /= -9 * [1111,2222,3333,4444])) stop 5 + endif + var4%y%y%x%x = 6 * var4%y%y%x%x + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + var4%y%y%x2(1)%x = 6 * var4%y%y%x2(1)%x + var4%y%y%x2(2)%x = 6 * var4%y%y%x2(2)%x + endif + var4%y2(1)%y%x%x = 6 * var4%y2(1)%y%x%x + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + var4%y2(1)%y%x2(1)%x = 6 * var4%y2(1)%y%x2(1)%x + var4%y2(1)%y%x2(2)%x = 6 * var4%y2(1)%y%x2(2)%x + endif + var4%y2(2)%y%x%x = 6 * var4%y2(2)%y%x%x + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + var4%y2(2)%y%x2(1)%x = 6 * var4%y2(2)%y%x2(1)%x + var4%y2(2)%y%x2(2)%x = 6 * var4%y2(2)%y%x2(2)%x + endif + !$omp end target + +else if (case == 5) then + ! Use target enter/exit data + target with explicit map + !$omp target enter data map(to: var1) + !$omp target enter data map(to: var1a) + !$omp target enter data map(to: var2) ! { dg-warning "35: Mapping of polymorphic list item 'var2' is unspecified behavior \\\[-Wopenmp\\\]" } + !$omp target enter data map(to: var3) ! { dg-warning "35: Mapping of polymorphic list item 'var3->x' is unspecified behavior \\\[-Wopenmp\\\]" } + !$omp target enter data map(to: var4) ! { dg-warning "35: Mapping of polymorphic list item 'var4\.\[0-9\]+->y->y\.x' is unspecified behavior \\\[-Wopenmp\\\]" } + + !$omp target map(to: var1) + if (any (var1%x /= [1,2,3,4])) stop 1 + var1%x = 2 * var1%x + !$omp end target + + !$omp target map(to: var1a) + if (any (var1a%x /= [-1,-2,-3,-4])) stop 2 + var1a%x = 3 * var1a%x + !$omp end target + + !$omp target map(to: var2) ! { dg-warning "24: Mapping of polymorphic list item 'var2' is unspecified behavior \\\[-Wopenmp\\\]" } + if (any (var2%x /= [11,22,33,44])) stop 3 + var2%x = 4 * var2%x + !$omp end target + + !$omp target map(to: var3) ! { dg-warning "24: Mapping of polymorphic list item 'var3->x' is unspecified behavior \\\[-Wopenmp\\\]" } + if (any (var3%x%x /= [111,222,333,444])) stop 4 + var3%x%x = 5 * var3%x%x + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + if (any (var3%x2(1)%x /= 2*[111,222,333,444])) stop 4 + if (any (var3%x2(2)%x /= 3*[111,222,333,444])) stop 4 + var3%x2(1)%x = 5 * var3%x2(1)%x + var3%x2(2)%x = 5 * var3%x2(2)%x + endif + !$omp end target + + !$omp target map(to: var4) ! { dg-warning "24: Mapping of polymorphic list item 'var4\.\[0-9\]+->y->y\.x' is unspecified behavior \\\[-Wopenmp\\\]" } + if (any (var4%y%y%x%x /= -1 * [1111,2222,3333,4444])) stop 5 + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + if (any (var4%y%y%x2(1)%x /= -2 * [1111,2222,3333,4444])) stop 5 + if (any (var4%y%y%x2(2)%x /= -3 * [1111,2222,3333,4444])) stop 5 + endif + if (any (var4%y2(1)%y%x%x /= -4 * [1111,2222,3333,4444])) stop 5 + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + if (any (var4%y2(1)%y%x2(1)%x /= -5 * [1111,2222,3333,4444])) stop 5 + if (any (var4%y2(1)%y%x2(2)%x /= -6 * [1111,2222,3333,4444])) stop 5 + endif + if (any (var4%y2(2)%y%x%x /= -7 * [1111,2222,3333,4444])) stop 5 + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + if (any (var4%y2(2)%y%x2(1)%x /= -8 * [1111,2222,3333,4444])) stop 5 + if (any (var4%y2(2)%y%x2(2)%x /= -9 * [1111,2222,3333,4444])) stop 5 + endif + var4%y%y%x%x = 6 * var4%y%y%x%x + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + var4%y%y%x2(1)%x = 6 * var4%y%y%x2(1)%x + var4%y%y%x2(2)%x = 6 * var4%y%y%x2(2)%x + endif + var4%y2(1)%y%x%x = 6 * var4%y2(1)%y%x%x + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + var4%y2(1)%y%x2(1)%x = 6 * var4%y2(1)%y%x2(1)%x + var4%y2(1)%y%x2(2)%x = 6 * var4%y2(1)%y%x2(2)%x + endif + var4%y2(2)%y%x%x = 6 * var4%y2(2)%y%x%x + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + var4%y2(2)%y%x2(1)%x = 6 * var4%y2(2)%y%x2(1)%x + var4%y2(2)%y%x2(2)%x = 6 * var4%y2(2)%y%x2(2)%x + endif + !$omp end target + + !$omp target exit data map(from: var1) + !$omp target exit data map(from: var1a) + !$omp target exit data map(from: var2) ! { dg-warning "36: Mapping of polymorphic list item 'var2' is unspecified behavior \\\[-Wopenmp\\\]" } + !$omp target exit data map(from: var3) ! { dg-warning "36: Mapping of polymorphic list item 'var3->x' is unspecified behavior \\\[-Wopenmp\\\]" } + !$omp target exit data map(from: var4) ! { dg-warning "36: Mapping of polymorphic list item 'var4\.\[0-9\]+->y->y\.x' is unspecified behavior \\\[-Wopenmp\\\]" } + +else if (case == 6) then + ! Use target enter/exit data + target with implicit map + + !$omp target enter data map(to: var1) + !$omp target enter data map(to: var1a) + !$omp target enter data map(to: var2) ! { dg-warning "35: Mapping of polymorphic list item 'var2' is unspecified behavior \\\[-Wopenmp\\\]" } + !$omp target enter data map(to: var3) ! { dg-warning "35: Mapping of polymorphic list item 'var3->x' is unspecified behavior \\\[-Wopenmp\\\]" } + !$omp target enter data map(to: var4) ! { dg-warning "35: Mapping of polymorphic list item 'var4\.\[0-9\]+->y->y\.x' is unspecified behavior \\\[-Wopenmp\\\]" } + + !$omp target + if (any (var1%x /= [1,2,3,4])) stop 1 + var1%x = 2 * var1%x + !$omp end target + + !$omp target + if (any (var1a%x /= [-1,-2,-3,-4])) stop 2 + var1a%x = 3 * var1a%x + !$omp end target + + !$omp target ! { dg-warning "Mapping of polymorphic list item 'var2' is unspecified behavior \\\[-Wopenmp\\\]" } + if (any (var2%x /= [11,22,33,44])) stop 3 + var2%x = 4 * var2%x + !$omp end target + + !$omp target ! { dg-warning "Mapping of polymorphic list item 'var3->x' is unspecified behavior \\\[-Wopenmp\\\]" } + if (any (var3%x%x /= [111,222,333,444])) stop 4 + var3%x%x = 5 * var3%x%x + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + if (any (var3%x2(1)%x /= 2*[111,222,333,444])) stop 4 + if (any (var3%x2(2)%x /= 3*[111,222,333,444])) stop 4 + var3%x2(1)%x = 5 * var3%x2(1)%x + var3%x2(2)%x = 5 * var3%x2(2)%x + endif + !$omp end target + + !$omp target ! { dg-warning "Mapping of polymorphic list item 'var4\.\[0-9\]+->y->y\.x' is unspecified behavior \\\[-Wopenmp\\\]" } + if (any (var4%y%y%x%x /= -1 * [1111,2222,3333,4444])) stop 5 + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + if (any (var4%y%y%x2(1)%x /= -2 * [1111,2222,3333,4444])) stop 5 + if (any (var4%y%y%x2(2)%x /= -3 * [1111,2222,3333,4444])) stop 5 + endif + if (any (var4%y2(1)%y%x%x /= -4 * [1111,2222,3333,4444])) stop 5 + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + if (any (var4%y2(1)%y%x2(1)%x /= -5 * [1111,2222,3333,4444])) stop 5 + if (any (var4%y2(1)%y%x2(2)%x /= -6 * [1111,2222,3333,4444])) stop 5 + endif + if (any (var4%y2(2)%y%x%x /= -7 * [1111,2222,3333,4444])) stop 5 + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + if (any (var4%y2(2)%y%x2(1)%x /= -8 * [1111,2222,3333,4444])) stop 5 + if (any (var4%y2(2)%y%x2(2)%x /= -9 * [1111,2222,3333,4444])) stop 5 + endif + var4%y%y%x%x = 6 * var4%y%y%x%x + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + var4%y%y%x2(1)%x = 6 * var4%y%y%x2(1)%x + var4%y%y%x2(2)%x = 6 * var4%y%y%x2(2)%x + endif + var4%y2(1)%y%x%x = 6 * var4%y2(1)%y%x%x + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + var4%y2(1)%y%x2(1)%x = 6 * var4%y2(1)%y%x2(1)%x + var4%y2(1)%y%x2(2)%x = 6 * var4%y2(1)%y%x2(2)%x + endif + var4%y2(2)%y%x%x = 6 * var4%y2(2)%y%x%x + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + var4%y2(2)%y%x2(1)%x = 6 * var4%y2(2)%y%x2(1)%x + var4%y2(2)%y%x2(2)%x = 6 * var4%y2(2)%y%x2(2)%x + endif + !$omp end target + + !$omp target exit data map(from: var1) + !$omp target exit data map(from: var1a) + !$omp target exit data map(from: var2) ! { dg-warning "36: Mapping of polymorphic list item 'var2' is unspecified behavior \\\[-Wopenmp\\\]" } + !$omp target exit data map(from: var3) ! { dg-warning "36: Mapping of polymorphic list item 'var3->x' is unspecified behavior \\\[-Wopenmp\\\]" } + !$omp target exit data map(from: var4) ! { dg-warning "36: Mapping of polymorphic list item 'var4\.\[0-9\]+->y->y\.x' is unspecified behavior \\\[-Wopenmp\\\]" } + +else + error stop +end if + +if ((case /= 2 .and. case /= 4) .or. is_shared_mem) then + ! The target update should have been active, check for the updated values + if (any (var1%x /= 2 * [1,2,3,4])) stop 11 + if (any (var1a%x /= 3 * [-1,-2,-3,-4])) stop 22 + if (any (var2%x /= 4 * [11,22,33,44])) stop 33 + + if (any (var3%x%x /= 5 * [111,222,333,444])) stop 44 + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + if (any (var3%x2(1)%x /= 2 * 5 * [111,222,333,444])) stop 44 + if (any (var3%x2(2)%x /= 3 * 5 * [111,222,333,444])) stop 44 + endif + + if (any (var4%y%y%x%x /= -1 * 6 * [1111,2222,3333,4444])) stop 55 + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + if (any (var4%y%y%x2(1)%x /= -2 * 6 * [1111,2222,3333,4444])) stop 55 + if (any (var4%y%y%x2(2)%x /= -3 * 6 * [1111,2222,3333,4444])) stop 55 + endif + if (any (var4%y2(1)%y%x%x /= -4 * 6 * [1111,2222,3333,4444])) stop 55 + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + if (any (var4%y2(1)%y%x2(1)%x /= -5 * 6 * [1111,2222,3333,4444])) stop 55 + if (any (var4%y2(1)%y%x2(2)%x /= -6 * 6 * [1111,2222,3333,4444])) stop 55 + endif + if (any (var4%y2(2)%y%x%x /= -7 * 6 * [1111,2222,3333,4444])) stop 55 + if (is_shared_mem) then ! For stride data, this accesses the host's _vtab + if (any (var4%y2(2)%y%x2(1)%x /= -8 * 6 * [1111,2222,3333,4444])) stop 55 + if (any (var4%y2(2)%y%x2(2)%x /= -9 * 6 * [1111,2222,3333,4444])) stop 55 + endif +else + ! The old host values should still be there as 'to:' created a device copy + if (any (var1%x /= [1,2,3,4])) stop 12 + if (any (var1a%x /= [-1,-2,-3,-4])) stop 22 + if (any (var2%x /= [11,22,33,44])) stop 33 + + if (any (var3%x%x /= [111,222,333,444])) stop 44 + ! .not. is_shared_mem: + ! if (any (var3%x2(1)%x /= 2*[111,222,333,444])) stop 44 + ! if (any (var3%x2(2)%x /= 3*[111,222,333,444])) stop 44 + + if (any (var4%y%y%x%x /= -1 * [1111,2222,3333,4444])) stop 55 + if (any (var4%y%y%x2(1)%x /= -2 * [1111,2222,3333,4444])) stop 55 + if (any (var4%y%y%x2(2)%x /= -3 * [1111,2222,3333,4444])) stop 55 + if (any (var4%y2(1)%y%x%x /= -4 * [1111,2222,3333,4444])) stop 55 + ! .not. is_shared_mem: + !if (any (var4%y2(1)%y%x2(1)%x /= -5 * [1111,2222,3333,4444])) stop 55 + !if (any (var4%y2(1)%y%x2(2)%x /= -6 * [1111,2222,3333,4444])) stop 55 + if (any (var4%y2(2)%y%x%x /= -7 * [1111,2222,3333,4444])) stop 55 + ! .not. is_shared_mem: + !if (any (var4%y2(2)%y%x2(1)%x /= -8 * [1111,2222,3333,4444])) stop 55 + !if (any (var4%y2(2)%y%x2(2)%x /= -9 * [1111,2222,3333,4444])) stop 55 +end if +if (case_var(100) /= 0) stop 123 +end subroutine test + +program main + use omp_lib + implicit none(type, external) + interface + subroutine test(case) + integer, value :: case + end + end interface + integer :: dev + call run_it(omp_get_default_device()) + do dev = 0, omp_get_num_devices() + call run_it(dev) + end do + call run_it(omp_initial_device) +! print *, 'all done' +contains +subroutine run_it(dev) + integer, value :: dev +! print *, 'DEVICE', dev + call omp_set_default_device(dev) + call test(1) + call test(2) + call test(3) + call test(4) + call test(5) + call test(6) +end +end diff --git a/libgomp/testsuite/libgomp.fortran/target-enter-data-8.f90 b/libgomp/testsuite/libgomp.fortran/target-enter-data-8.f90 new file mode 100644 index 0000000..c6d671c --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-enter-data-8.f90 @@ -0,0 +1,532 @@ +! { dg-additional-options "-cpp" } + +! FIXME: Some tests do not work yet. Those are for now in '#if 0' + +! Check that 'map(alloc:' properly works with +! - deferred-length character strings +! - arrays with array descriptors +! For those, the array descriptor / string length must be mapped with 'to:' + +program main +implicit none + +type t + integer :: ic(2:5), ic2 + character(len=11) :: ccstr(3:4), ccstr2 + character(len=11,kind=4) :: cc4str(3:7), cc4str2 + integer, pointer :: pc(:), pc2 + character(len=:), pointer :: pcstr(:), pcstr2 + character(len=:,kind=4), pointer :: pc4str(:), pc4str2 +end type t + +type(t) :: dt + +integer :: ii(5), ii2 +character(len=11) :: clstr(-1:1), clstr2 +character(len=11,kind=4) :: cl4str(0:3), cl4str2 +integer, pointer :: ip(:), ip2 +integer, allocatable :: ia(:), ia2 +character(len=:), pointer :: pstr(:), pstr2 +character(len=:), allocatable :: astr(:), astr2 +character(len=:,kind=4), pointer :: p4str(:), p4str2 +character(len=:,kind=4), allocatable :: a4str(:), a4str2 + + +allocate(dt%pc(5), dt%pc2) +allocate(character(len=2) :: dt%pcstr(2)) +allocate(character(len=4) :: dt%pcstr2) + +allocate(character(len=3,kind=4) :: dt%pc4str(2:3)) +allocate(character(len=5,kind=4) :: dt%pc4str2) + +allocate(ip(5), ip2, ia(8), ia2) +allocate(character(len=2) :: pstr(-2:0)) +allocate(character(len=4) :: pstr2) +allocate(character(len=6) :: astr(3:5)) +allocate(character(len=8) :: astr2) + +allocate(character(len=3,kind=4) :: p4str(2:4)) +allocate(character(len=5,kind=4) :: p4str2) +allocate(character(len=7,kind=4) :: a4str(-2:3)) +allocate(character(len=9,kind=4) :: a4str2) + + +! integer :: ic(2:5), ic2 + +!$omp target enter data map(alloc: dt%ic) +!$omp target map(alloc: dt%ic) + if (size(dt%ic) /= 4) error stop + if (lbound(dt%ic, 1) /= 2) error stop + if (ubound(dt%ic, 1) /= 5) error stop + dt%ic = [22, 33, 44, 55] +!$omp end target +!$omp target exit data map(from: dt%ic) +if (size(dt%ic) /= 4) error stop +if (lbound(dt%ic, 1) /= 2) error stop +if (ubound(dt%ic, 1) /= 5) error stop +if (any (dt%ic /= [22, 33, 44, 55])) error stop + +!$omp target enter data map(alloc: dt%ic2) +!$omp target map(alloc: dt%ic2) + dt%ic2 = 42 +!$omp end target +!$omp target exit data map(from: dt%ic2) +if (dt%ic2 /= 42) error stop + + +! character(len=11) :: ccstr(3:4), ccstr2 + +!$omp target enter data map(alloc: dt%ccstr) +!$omp target map(alloc: dt%ccstr) + if (len(dt%ccstr) /= 11) error stop + if (size(dt%ccstr) /= 2) error stop + if (lbound(dt%ccstr, 1) /= 3) error stop + if (ubound(dt%ccstr, 1) /= 4) error stop + dt%ccstr = ["12345678901", "abcdefghijk"] +!$omp end target +!$omp target exit data map(from: dt%ccstr) +if (len(dt%ccstr) /= 11) error stop +if (size(dt%ccstr) /= 2) error stop +if (lbound(dt%ccstr, 1) /= 3) error stop +if (ubound(dt%ccstr, 1) /= 4) error stop +if (any (dt%ccstr /= ["12345678901", "abcdefghijk"])) error stop + +!$omp target enter data map(alloc: dt%ccstr2) +!$omp target map(alloc: dt%ccstr2) + if (len(dt%ccstr2) /= 11) error stop + dt%ccstr2 = "ABCDEFGHIJK" +!$omp end target +!$omp target exit data map(from: dt%ccstr2) +if (len(dt%ccstr2) /= 11) error stop +if (dt%ccstr2 /= "ABCDEFGHIJK") error stop + + +! character(len=11,kind=4) :: cc4str(3:7), cc4str2 + +#if 0 +! Value check fails +!$omp target map(alloc: dt%cc4str) + if (len(dt%cc4str) /= 11) error stop + if (size(dt%cc4str) /= 5) error stop + if (lbound(dt%cc4str, 1) /= 3) error stop + if (ubound(dt%cc4str, 1) /= 7) error stop + dt%cc4str = [4_"12345678901", 4_"abcdefghijk", & + 4_"qerftcea6ds", 4_"a1f9g37ga4.", & + 4_"45ngwj56sj2"] +!$omp end target +!$omp target exit data map(from: dt%cc4str) +if (len(dt%cc4str) /= 11) error stop +if (size(dt%cc4str) /= 5) error stop +if (lbound(dt%cc4str, 1) /= 3) error stop +if (ubound(dt%cc4str, 1) /= 7) error stop +if (dt%cc4str(3) /= 4_"12345678901") error stop +if (dt%cc4str(4) /= 4_"abcdefghijk") error stop +if (dt%cc4str(5) /= 4_"qerftcea6ds") error stop +if (dt%cc4str(6) /= 4_"a1f9g37ga4.") error stop +if (dt%cc4str(7) /= 4_"45ngwj56sj2") error stop +#endif + +!$omp target enter data map(alloc: dt%cc4str2) +!$omp target map(alloc: dt%cc4str2) + if (len(dt%cc4str2) /= 11) error stop + dt%cc4str2 = 4_"ABCDEFGHIJK" +!$omp end target +!$omp target exit data map(from: dt%cc4str2) +if (len(dt%cc4str2) /= 11) error stop +if (dt%cc4str2 /= 4_"ABCDEFGHIJK") error stop + + +! integer, pointer :: pc(:), pc2 +! allocate(dt%pc(5), dt%pc2) + +!$omp target enter data map(alloc: dt%pc) +!$omp target map(alloc: dt%pc) + if (.not. associated(dt%pc)) error stop + if (size(dt%pc) /= 5) error stop + if (lbound(dt%pc, 1) /= 1) error stop + if (ubound(dt%pc, 1) /= 5) error stop + dt%pc = [11, 22, 33, 44, 55] +!$omp end target +!$omp target exit data map(from: dt%pc) +if (.not. associated(dt%pc)) error stop +if (size(dt%pc) /= 5) error stop +if (lbound(dt%pc, 1) /= 1) error stop +if (ubound(dt%pc, 1) /= 5) error stop +if (any (dt%pc /= [11, 22, 33, 44, 55])) error stop + +!$omp target enter data map(alloc: dt%pc2) +!$omp target map(alloc: dt%pc2) + if (.not. associated(dt%pc2)) error stop + dt%pc2 = 99 +!$omp end target +!$omp target exit data map(from: dt%pc2) +if (dt%pc2 /= 99) error stop +if (.not. associated(dt%pc2)) error stop + + +! character(len=:), pointer :: pcstr(:), pcstr2 +! allocate(character(len=2) :: dt%pcstr(2)) +! allocate(character(len=4) :: dt%pcstr2) + +!$omp target enter data map(alloc: dt%pcstr) +!$omp target map(alloc: dt%pcstr) + if (.not. associated(dt%pcstr)) error stop + if (len(dt%pcstr) /= 2) error stop + if (size(dt%pcstr) /= 2) error stop + if (lbound(dt%pcstr, 1) /= 1) error stop + if (ubound(dt%pcstr, 1) /= 2) error stop + dt%pcstr = ["01", "jk"] +!$omp end target +!$omp target exit data map(from: dt%pcstr) +if (.not. associated(dt%pcstr)) error stop +if (len(dt%pcstr) /= 2) error stop +if (size(dt%pcstr) /= 2) error stop +if (lbound(dt%pcstr, 1) /= 1) error stop +if (ubound(dt%pcstr, 1) /= 2) error stop +if (any (dt%pcstr /= ["01", "jk"])) error stop + + +!$omp target enter data map(alloc: dt%pcstr2) +!$omp target map(alloc: dt%pcstr2) + if (.not. associated(dt%pcstr2)) error stop + if (len(dt%pcstr2) /= 4) error stop + dt%pcstr2 = "HIJK" +!$omp end target +!$omp target exit data map(from: dt%pcstr2) +if (.not. associated(dt%pcstr2)) error stop +if (len(dt%pcstr2) /= 4) error stop +if (dt%pcstr2 /= "HIJK") error stop + + +! character(len=:,kind=4), pointer :: pc4str(:), pc4str2 +! allocate(character(len=3,kind=4) :: dt%pc4str(2:3)) +! allocate(character(len=5,kind=4) :: dt%pc4str2) + +!$omp target enter data map(alloc: dt%pc4str) +!$omp target map(alloc: dt%pc4str) + if (.not. associated(dt%pc4str)) error stop + if (len(dt%pc4str) /= 3) error stop + if (size(dt%pc4str) /= 2) error stop + if (lbound(dt%pc4str, 1) /= 2) error stop + if (ubound(dt%pc4str, 1) /= 3) error stop + dt%pc4str = [4_"456", 4_"tzu"] +!$omp end target +!$omp target exit data map(from: dt%pc4str) +if (.not. associated(dt%pc4str)) error stop +if (len(dt%pc4str) /= 3) error stop +if (size(dt%pc4str) /= 2) error stop +if (lbound(dt%pc4str, 1) /= 2) error stop +if (ubound(dt%pc4str, 1) /= 3) error stop +if (dt%pc4str(2) /= 4_"456") error stop +if (dt%pc4str(3) /= 4_"tzu") error stop + +!$omp target enter data map(alloc: dt%pc4str2) +!$omp target map(alloc: dt%pc4str2) + if (.not. associated(dt%pc4str2)) error stop + if (len(dt%pc4str2) /= 5) error stop + dt%pc4str2 = 4_"98765" +!$omp end target +!$omp target exit data map(from: dt%pc4str2) +if (.not. associated(dt%pc4str2)) error stop +if (len(dt%pc4str2) /= 5) error stop +if (dt%pc4str2 /= 4_"98765") error stop + + +! integer :: ii(5), ii2 + +!$omp target enter data map(alloc: ii) +!$omp target map(alloc: ii) + if (size(ii) /= 5) error stop + if (lbound(ii, 1) /= 1) error stop + if (ubound(ii, 1) /= 5) error stop + ii = [-1, -2, -3, -4, -5] +!$omp end target +!$omp target exit data map(from: ii) +if (size(ii) /= 5) error stop +if (lbound(ii, 1) /= 1) error stop +if (ubound(ii, 1) /= 5) error stop +if (any (ii /= [-1, -2, -3, -4, -5])) error stop + +!$omp target enter data map(alloc: ii2) +!$omp target map(alloc: ii2) + ii2 = -410 +!$omp end target +!$omp target exit data map(from: ii2) +if (ii2 /= -410) error stop + + +! character(len=11) :: clstr(-1:1), clstr2 + +!$omp target enter data map(alloc: clstr) +!$omp target map(alloc: clstr) + if (len(clstr) /= 11) error stop + if (size(clstr) /= 3) error stop + if (lbound(clstr, 1) /= -1) error stop + if (ubound(clstr, 1) /= 1) error stop + clstr = ["12345678901", "abcdefghijk", "ABCDEFGHIJK"] +!$omp end target +!$omp target exit data map(from: clstr) +if (len(clstr) /= 11) error stop +if (size(clstr) /= 3) error stop +if (lbound(clstr, 1) /= -1) error stop +if (ubound(clstr, 1) /= 1) error stop +if (any (clstr /= ["12345678901", "abcdefghijk", "ABCDEFGHIJK"])) error stop + +!$omp target enter data map(alloc: clstr2) +!$omp target map(alloc: clstr2) + if (len(clstr2) /= 11) error stop + clstr2 = "ABCDEFghijk" +!$omp end target +!$omp target exit data map(from: clstr2) +if (len(clstr2) /= 11) error stop +if (clstr2 /= "ABCDEFghijk") error stop + + +! character(len=11,kind=4) :: cl4str(0:3), cl4str2 + +!$omp target enter data map(alloc: cl4str) +!$omp target map(alloc: cl4str) + if (len(cl4str) /= 11) error stop + if (size(cl4str) /= 4) error stop + if (lbound(cl4str, 1) /= 0) error stop + if (ubound(cl4str, 1) /= 3) error stop + cl4str = [4_"12345678901", 4_"abcdefghijk", & + 4_"qerftcea6ds", 4_"a1f9g37ga4."] +!$omp end target +!$omp target exit data map(from: cl4str) +if (len(cl4str) /= 11) error stop +if (size(cl4str) /= 4) error stop +if (lbound(cl4str, 1) /= 0) error stop +if (ubound(cl4str, 1) /= 3) error stop +if (cl4str(0) /= 4_"12345678901") error stop +if (cl4str(1) /= 4_"abcdefghijk") error stop +if (cl4str(2) /= 4_"qerftcea6ds") error stop +if (cl4str(3) /= 4_"a1f9g37ga4.") error stop + +!$omp target enter data map(alloc: cl4str2) +!$omp target map(alloc: cl4str2) + if (len(cl4str2) /= 11) error stop + cl4str2 = 4_"ABCDEFGHIJK" +!$omp end target +!$omp target exit data map(from: cl4str2) +if (len(cl4str2) /= 11) error stop +if (cl4str2 /= 4_"ABCDEFGHIJK") error stop + + +! allocate(ip(5), ip2, ia(8), ia2) + +!$omp target enter data map(alloc: ip) +!$omp target map(alloc: ip) + if (.not. associated(ip)) error stop + if (size(ip) /= 5) error stop + if (lbound(ip, 1) /= 1) error stop + if (ubound(ip, 1) /= 5) error stop + ip = [11, 22, 33, 44, 55] +!$omp end target +!$omp target exit data map(from: ip) +if (.not. associated(ip)) error stop +if (size(ip) /= 5) error stop +if (lbound(ip, 1) /= 1) error stop +if (ubound(ip, 1) /= 5) error stop +if (any (ip /= [11, 22, 33, 44, 55])) error stop + +!$omp target enter data map(alloc: ip2) +!$omp target map(alloc: ip2) + if (.not. associated(ip2)) error stop + ip2 = 99 +!$omp end target +!$omp target exit data map(from: ip2) +if (ip2 /= 99) error stop +if (.not. associated(ip2)) error stop + + +! allocate(ip(5), ip2, ia(8), ia2) + +!$omp target enter data map(alloc: ia) +!$omp target map(alloc: ia) + if (.not. allocated(ia)) error stop + if (size(ia) /= 8) error stop + if (lbound(ia, 1) /= 1) error stop + if (ubound(ia, 1) /= 8) error stop + ia = [1,2,3,4,5,6,7,8] +!$omp end target +!$omp target exit data map(from: ia) +if (.not. allocated(ia)) error stop +if (size(ia) /= 8) error stop +if (lbound(ia, 1) /= 1) error stop +if (ubound(ia, 1) /= 8) error stop +if (any (ia /= [1,2,3,4,5,6,7,8])) error stop + +!$omp target enter data map(alloc: ia2) +!$omp target map(alloc: ia2) + if (.not. allocated(ia2)) error stop + ia2 = 102 +!$omp end target +!$omp target exit data map(from: ia2) +if (ia2 /= 102) error stop +if (.not. allocated(ia2)) error stop + + +! character(len=:), pointer :: pstr(:), pstr2 +! allocate(character(len=2) :: pstr(-2:0)) +! allocate(character(len=4) :: pstr2) + +!$omp target enter data map(alloc: pstr) +!$omp target map(alloc: pstr) + if (.not. associated(pstr)) error stop + if (len(pstr) /= 2) error stop + if (size(pstr) /= 3) error stop + if (lbound(pstr, 1) /= -2) error stop + if (ubound(pstr, 1) /= 0) error stop + pstr = ["01", "jk", "aq"] +!$omp end target +!$omp target exit data map(from: pstr) +if (.not. associated(pstr)) error stop +if (len(pstr) /= 2) error stop +if (size(pstr) /= 3) error stop +if (lbound(pstr, 1) /= -2) error stop +if (ubound(pstr, 1) /= 0) error stop +if (any (pstr /= ["01", "jk", "aq"])) error stop + +!$omp target enter data map(alloc: pstr2) +!$omp target map(alloc: pstr2) + if (.not. associated(pstr2)) error stop + if (len(pstr2) /= 4) error stop + pstr2 = "HIJK" +!$omp end target +!$omp target exit data map(from: pstr2) +if (.not. associated(pstr2)) error stop +if (len(pstr2) /= 4) error stop +if (pstr2 /= "HIJK") error stop + + +! character(len=:), allocatable :: astr(:), astr2 +! allocate(character(len=6) :: astr(3:5)) +! allocate(character(len=8) :: astr2) + + +!$omp target enter data map(alloc: astr) +!$omp target map(alloc: astr) + if (.not. allocated(astr)) error stop + if (len(astr) /= 6) error stop + if (size(astr) /= 3) error stop + if (lbound(astr, 1) /= 3) error stop + if (ubound(astr, 1) /= 5) error stop + astr = ["01db45", "jk$D%S", "zutg47"] +!$omp end target +!$omp target exit data map(from: astr) +if (.not. allocated(astr)) error stop +if (len(astr) /= 6) error stop +if (size(astr) /= 3) error stop +if (lbound(astr, 1) /= 3) error stop +if (ubound(astr, 1) /= 5) error stop +if (any (astr /= ["01db45", "jk$D%S", "zutg47"])) error stop + + +!$omp target enter data map(alloc: astr2) +!$omp target map(alloc: astr2) + if (.not. allocated(astr2)) error stop + if (len(astr2) /= 8) error stop + astr2 = "HIJKhijk" +!$omp end target +!$omp target exit data map(from: astr2) +if (.not. allocated(astr2)) error stop +if (len(astr2) /= 8) error stop +if (astr2 /= "HIJKhijk") error stop + + +! character(len=:,kind=4), pointer :: p4str(:), p4str2 +! allocate(character(len=3,kind=4) :: p4str(2:4)) +! allocate(character(len=5,kind=4) :: p4str2) + +! FAILS with value check + +!$omp target enter data map(alloc: p4str) +!$omp target map(alloc: p4str) + if (.not. associated(p4str)) error stop + if (len(p4str) /= 3) error stop + if (size(p4str) /= 3) error stop + if (lbound(p4str, 1) /= 2) error stop + if (ubound(p4str, 1) /= 4) error stop + p4str(:) = [4_"f85", 4_"8af", 4_"A%F"] +!$omp end target +!$omp target exit data map(from: p4str) +if (.not. associated(p4str)) error stop +if (len(p4str) /= 3) error stop +if (size(p4str) /= 3) error stop +if (lbound(p4str, 1) /= 2) error stop +if (ubound(p4str, 1) /= 4) error stop +if (p4str(2) /= 4_"f85") error stop +if (p4str(3) /= 4_"8af") error stop +if (p4str(4) /= 4_"A%F") error stop + +!$omp target enter data map(alloc: p4str2) +!$omp target map(alloc: p4str2) + if (.not. associated(p4str2)) error stop + if (len(p4str2) /= 5) error stop + p4str2 = 4_"9875a" +!$omp end target +!$omp target exit data map(from: p4str2) +if (.not. associated(p4str2)) error stop +if (len(p4str2) /= 5) error stop +if (p4str2 /= 4_"9875a") error stop + + +! character(len=:,kind=4), allocatable :: a4str(:), a4str2 +! allocate(character(len=7,kind=4) :: a4str(-2:3)) +! allocate(character(len=9,kind=4) :: a4str2) + +!$omp target enter data map(alloc: a4str) +!$omp target map(alloc: a4str) + if (.not. allocated(a4str)) error stop + if (len(a4str) /= 7) error stop + if (size(a4str) /= 6) error stop + if (lbound(a4str, 1) /= -2) error stop + if (ubound(a4str, 1) /= 3) error stop + ! See PR fortran/107508 why '(:)' is required + a4str(:) = [4_"sf456aq", 4_"3dtzu24", 4_"_4fh7sm", 4_"=ff85s7", 4_"j=8af4d", 4_".,A%Fsz"] +!$omp end target +!$omp target exit data map(from: a4str) +if (.not. allocated(a4str)) error stop +if (len(a4str) /= 7) error stop +if (size(a4str) /= 6) error stop +if (lbound(a4str, 1) /= -2) error stop +if (ubound(a4str, 1) /= 3) error stop +if (a4str(-2) /= 4_"sf456aq") error stop +if (a4str(-1) /= 4_"3dtzu24") error stop +if (a4str(0) /= 4_"_4fh7sm") error stop +if (a4str(1) /= 4_"=ff85s7") error stop +if (a4str(2) /= 4_"j=8af4d") error stop +if (a4str(3) /= 4_".,A%Fsz") error stop + +!$omp target enter data map(alloc: a4str2) +!$omp target map(alloc: a4str2) + if (.not. allocated(a4str2)) error stop + if (len(a4str2) /= 9) error stop + a4str2 = 4_"98765a23d" +!$omp end target +!$omp target exit data map(from: a4str2) +if (.not. allocated(a4str2)) error stop +if (len(a4str2) /= 9) error stop +if (a4str2 /= 4_"98765a23d") error stop + + +deallocate(dt%pc, dt%pc2) +deallocate(dt%pcstr) +deallocate(dt%pcstr2) + +deallocate(dt%pc4str) +deallocate(dt%pc4str2) + +deallocate(ip, ip2, ia, ia2) +deallocate(pstr) +deallocate(pstr2) +deallocate(astr) +deallocate(astr2) + +deallocate(p4str) +deallocate(p4str2) +deallocate(a4str) +deallocate(a4str2) + +end diff --git a/libgomp/testsuite/libgomp.oacc-c++/exceptions-bad_cast-1.C b/libgomp/testsuite/libgomp.oacc-c++/exceptions-bad_cast-1.C new file mode 100644 index 0000000..0545601 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/exceptions-bad_cast-1.C @@ -0,0 +1,54 @@ +/* 'std::bad_cast' exception in OpenACC compute 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 } */ + +/* See also '../libgomp.c++/target-exceptions-bad_cast-1.C'. */ + +/* See also '../../../gcc/testsuite/g++.target/gcn/exceptions-bad_cast-1.C', + '../../../gcc/testsuite/g++.target/nvptx/exceptions-bad_cast-1.C'. */ + +#include <iostream> + +struct C1 +{ + virtual void f() + {} +}; + +struct C2 : C1 +{ +}; + +int main() +{ + std::cerr << "CheCKpOInT\n"; +#pragma omp target +#pragma acc serial + /* { dg-bogus {using 'vector_length \(32\)', ignoring 1} {} { target openacc_nvidia_accel_selected xfail *-*-* } .-1 } */ + { + C1 c1; + [[maybe_unused]] + C2 &c2 = dynamic_cast<C2 &>(c1); + /* 'std::bad_cast' is thrown. */ + } +} + +/* { 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 openacc_host_selected } } + 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.oacc-c++/exceptions-bad_cast-2-offload-sorry-GCN.C b/libgomp/testsuite/libgomp.oacc-c++/exceptions-bad_cast-2-offload-sorry-GCN.C new file mode 100644 index 0000000..8260966 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/exceptions-bad_cast-2-offload-sorry-GCN.C @@ -0,0 +1,18 @@ +/* 'std::bad_cast' exception in OpenACC compute 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 openacc_radeon_accel_selected } } */ +/* { 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 "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.oacc-c++/exceptions-bad_cast-2-offload-sorry-nvptx.C b/libgomp/testsuite/libgomp.oacc-c++/exceptions-bad_cast-2-offload-sorry-nvptx.C new file mode 100644 index 0000000..86d3f6c --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/exceptions-bad_cast-2-offload-sorry-nvptx.C @@ -0,0 +1,20 @@ +/* 'std::bad_cast' exception in OpenACC compute 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 openacc_nvidia_accel_selected } } */ +/* { 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 "exceptions-bad_cast-2.C" + +/* { dg-bogus {using 'vector_length \(32\)', ignoring 1} {} { target openacc_nvidia_accel_selected xfail *-*-* } 0 } */ + +/* { 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.oacc-c++/exceptions-bad_cast-2.C b/libgomp/testsuite/libgomp.oacc-c++/exceptions-bad_cast-2.C new file mode 100644 index 0000000..24399ef --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/exceptions-bad_cast-2.C @@ -0,0 +1,60 @@ +/* 'std::bad_cast' exception in OpenACC compute 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 openacc_nvidia_accel_selected xfail *-*-* } 0 } + { dg-excess-errors {'mkoffload' failure etc.} { xfail openacc_nvidia_accel_selected } } */ + +/* See also '../libgomp.c++/target-exceptions-bad_cast-2.C'. */ + +/* See also '../../../gcc/testsuite/g++.target/gcn/exceptions-bad_cast-2.C', + '../../../gcc/testsuite/g++.target/nvptx/exceptions-bad_cast-2.C'. */ + +#include <iostream> +#include <typeinfo> + +struct C1 +{ + virtual void f() + {} +}; + +struct C2 : C1 +{ +}; + +int main() +{ + std::cerr << "CheCKpOInT\n"; +#pragma omp target +#pragma acc serial + { + C1 c1; + try + { + [[maybe_unused]] + C2 &c2 = dynamic_cast<C2 &>(c1); + /* 'std::bad_cast' is thrown. */ + } + catch (const std::bad_cast &e) + { + __builtin_printf("caught '%s'\n", e.what()); + } + } +} + +/* { 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 openacc_host_selected } } + 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 {'std::bad_cast' exception} { ! openacc_host_selected } } */ diff --git a/libgomp/testsuite/libgomp.oacc-c++/exceptions-bad_cast-3.C b/libgomp/testsuite/libgomp.oacc-c++/exceptions-bad_cast-3.C new file mode 100644 index 0000000..4fa419f --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/exceptions-bad_cast-3.C @@ -0,0 +1,49 @@ +/* 'std::bad_cast' exception in OpenACC compute region, dead code. */ + +/* { dg-require-effective-target exceptions } + { dg-additional-options -fexceptions } */ +/* Wrong code for offloading execution. + { dg-skip-if PR119692 { ! openacc_host_selected } } + { dg-additional-options -fdump-tree-gimple } */ +/* { dg-additional-options -fdump-tree-optimized-raw } */ + +/* See also '../libgomp.c++/target-exceptions-bad_cast-3.C'. */ + +/* See also '../../../gcc/testsuite/g++.target/gcn/exceptions-bad_cast-3.C', + '../../../gcc/testsuite/g++.target/nvptx/exceptions-bad_cast-3.C'. */ + +/* For PR119692 workarounds. */ +#ifndef DEFAULT +# define DEFAULT +#endif + +struct C1 +{ + virtual void f() + {} +}; + +struct C2 : C1 +{ +}; + +int main() +{ +#pragma omp target DEFAULT +#pragma acc serial DEFAULT + { + C1 c1; + bool a = false; + asm volatile ("" : : "r" (&a) : "memory"); + if (a) + { + [[maybe_unused]] + C2 &c2 = dynamic_cast<C2 &>(c1); + /* 'std::bad_cast' is thrown. */ + } + } +} + +/* { dg-final { scan-tree-dump-not {(?n)#pragma omp target oacc_serial map\(tofrom:_ZTI2C2 \[len: [0-9]+\]\) map\(tofrom:_ZTI2C1 \[len: [0-9]+\]\) map\(tofrom:_ZTV2C1 \[len: [0-9]+\]\)$} gimple { xfail *-*-* } } } */ + +/* { dg-final { scan-tree-dump-times {gimple_call <__cxa_bad_cast, } 1 optimized } } */ diff --git a/libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-1.C b/libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-1.C new file mode 100644 index 0000000..f2ef751 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-1.C @@ -0,0 +1,46 @@ +/* 'throw' in OpenACC compute 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 { openacc_radeon_accel_selected && __OPTIMIZE__ } xfail *-*-* } 0 } + { dg-ice PR119737 { openacc_radeon_accel_selected && __OPTIMIZE__ } } + { dg-excess-errors {'mkoffload' failure etc.} { xfail { openacc_radeon_accel_selected && __OPTIMIZE__ } } } */ + +/* See also '../libgomp.c++/target-exceptions-throw-1.C'. */ + +/* See also '../../../gcc/testsuite/g++.target/gcn/exceptions-throw-1.C', + '../../../gcc/testsuite/g++.target/nvptx/exceptions-throw-1.C'. */ + +#include <iostream> + +class MyException +{ +}; + +int main() +{ + std::cerr << "CheCKpOInT\n"; +#pragma omp target +#pragma acc serial + /* { dg-bogus {using 'vector_length \(32\)', ignoring 1} {} { target openacc_nvidia_accel_selected xfail *-*-* } .-1 } */ + { + MyException e1; + throw e1; + } +} + +/* { 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 openacc_host_selected } } + For GCN, nvptx offload execution, we don't print anything, but just 'abort'. + + { dg-shouldfail {'MyException' exception} } */ diff --git a/libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-2-offload-sorry-GCN.C b/libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-2-offload-sorry-GCN.C new file mode 100644 index 0000000..40be837 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-2-offload-sorry-GCN.C @@ -0,0 +1,20 @@ +/* 'throw' in OpenACC compute 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 openacc_radeon_accel_selected } } */ +/* { 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 "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.oacc-c++/exceptions-throw-2-offload-sorry-nvptx.C b/libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-2-offload-sorry-nvptx.C new file mode 100644 index 0000000..9461455 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-2-offload-sorry-nvptx.C @@ -0,0 +1,22 @@ +/* 'throw' in OpenACC compute 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 openacc_nvidia_accel_selected } } */ +/* { 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 "exceptions-throw-2.C" + +/* { dg-bogus {using 'vector_length \(32\)', ignoring 1} {} { target openacc_nvidia_accel_selected xfail *-*-* } 0 } */ + +/* { 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.oacc-c++/exceptions-throw-2.C b/libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-2.C new file mode 100644 index 0000000..f6dc970 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-2.C @@ -0,0 +1,55 @@ +/* 'throw' in OpenACC compute 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 {undefined symbol: typeinfo name for MyException} PR119806 { target { openacc_radeon_accel_selected && { ! __OPTIMIZE__ } } xfail *-*-* } 0 } + { dg-excess-errors {'mkoffload' failure etc.} { xfail { openacc_radeon_accel_selected && { ! __OPTIMIZE__ } } } } */ +/* { dg-bogus {Size expression must be absolute\.} PR119737 { target { openacc_radeon_accel_selected && __OPTIMIZE__ } xfail *-*-* } 0 } + { dg-ice PR119737 { openacc_radeon_accel_selected && __OPTIMIZE__ } } + { dg-excess-errors {'mkoffload' failures etc.} { xfail { openacc_radeon_accel_selected && __OPTIMIZE__ } } } */ +/* { dg-bogus {Initial value type mismatch} PR119806 { target { openacc_nvidia_accel_selected && { ! __OPTIMIZE__ } } xfail *-*-* } 0 } + { dg-excess-errors {'mkoffload' failure etc.} { xfail { openacc_nvidia_accel_selected && { ! __OPTIMIZE__ } } } } */ + +/* See also '../libgomp.c++/target-exceptions-throw-2.C'. */ + +/* See also '../../../gcc/testsuite/g++.target/gcn/exceptions-throw-2.C', + '../../../gcc/testsuite/g++.target/nvptx/exceptions-throw-2.C'. */ + +#include <iostream> + +class MyException +{ +}; + +int main() +{ + std::cerr << "CheCKpOInT\n"; +#pragma omp target +#pragma acc serial + /* { dg-bogus {using 'vector_length \(32\)', ignoring 1} {} { target openacc_nvidia_accel_selected xfail *-*-* } .-1 } */ + { + try + { + MyException e1; + throw e1; + } + catch (const MyException &e) + { + __builtin_printf("caught '%s'\n", "MyException"); + } + } +} + +/* { 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 openacc_host_selected } } + 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} { ! openacc_host_selected } } */ diff --git a/libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-3.C b/libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-3.C new file mode 100644 index 0000000..74a62b3 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/exceptions-throw-3.C @@ -0,0 +1,43 @@ +/* 'throw' in OpenACC compute region, dead code. */ + +/* { dg-require-effective-target exceptions } + { dg-additional-options -fexceptions } */ +/* Wrong code for offloading execution. + { dg-skip-if PR119692 { ! openacc_host_selected } } + { dg-additional-options -fdump-tree-gimple } */ +/* { dg-additional-options -fdump-tree-optimized-raw } */ + +/* See also '../libgomp.c++/target-exceptions-throw-3.C'. */ + +/* See also '../../../gcc/testsuite/g++.target/gcn/exceptions-throw-3.C', + '../../../gcc/testsuite/g++.target/nvptx/exceptions-throw-3.C'. */ + +/* For PR119692 workarounds. */ +#ifndef DEFAULT +# define DEFAULT +#endif + +class MyException +{ +}; + +int main() +{ +#pragma omp target DEFAULT +#pragma acc serial DEFAULT + /* { dg-bogus {using 'vector_length \(32\)', ignoring 1} {} { target openacc_nvidia_accel_selected xfail *-*-* } .-1 } */ + { + bool a = false; + asm volatile ("" : : "r" (&a) : "memory"); + if (a) + { + MyException e1; + throw e1; + } + } +} + +/* { dg-final { scan-tree-dump-not {(?n)#pragma omp target oacc_serial map\(tofrom:_ZTI11MyException \[len: [0-9]+\]\)$} 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 } } */ diff --git a/libgomp/testsuite/libgomp.oacc-c++/pr119692-1-1.C b/libgomp/testsuite/libgomp.oacc-c++/pr119692-1-1.C new file mode 100644 index 0000000..5c3e037 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/pr119692-1-1.C @@ -0,0 +1,42 @@ +/* PR119692 "C++ 'typeinfo', 'vtable' vs. OpenACC, OpenMP 'target' offloading" */ + +/* { dg-additional-options -UDEFAULT } + Wrong code for offloading execution. + { dg-skip-if PR119692 { ! openacc_host_selected } } */ +/* { dg-additional-options -fdump-tree-gimple } */ + +/* See also '../libgomp.c++/pr119692-1-1.C'. */ + +/* See also '../../../gcc/testsuite/g++.target/gcn/pr119692-1-1.C', + '../../../gcc/testsuite/g++.target/nvptx/pr119692-1-1.C'. */ + +#ifndef DEFAULT +# define DEFAULT +#endif + +struct C1 +{ + virtual void f() + {} +}; + +struct C2 : C1 +{ +}; + +int main() +{ +#pragma omp target DEFAULT +#pragma acc serial DEFAULT + /* { dg-bogus {using 'vector_length \(32\)', ignoring 1} {} { target openacc_nvidia_accel_selected xfail *-*-* } .-1 } */ + { + C1 c1; + C1 *c1p = &c1; + asm volatile ("" : : "r" (&c1p) : "memory"); + C2 *c2 = dynamic_cast<C2 *>(c1p); + if (c2) + __builtin_abort(); + } +} + +/* { dg-final { scan-tree-dump-not {(?n)#pragma omp target oacc_serial map\(tofrom:_ZTI2C2 \[len: [0-9]+\]\) map\(tofrom:_ZTI2C1 \[len: [0-9]+\]\) map\(tofrom:_ZTV2C1 \[len: [0-9]+\]\)$} gimple { xfail *-*-* } } } */ diff --git a/libgomp/testsuite/libgomp.oacc-c++/pr119692-1-2.C b/libgomp/testsuite/libgomp.oacc-c++/pr119692-1-2.C new file mode 100644 index 0000000..207b183 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/pr119692-1-2.C @@ -0,0 +1,12 @@ +/* PR119692 "C++ 'typeinfo', 'vtable' vs. OpenACC, OpenMP 'target' offloading" */ + +/* { dg-additional-options -DDEFAULT=default(none) } + Wrong code for offloading execution. + { dg-skip-if PR119692 { ! openacc_host_selected } } */ +/* { dg-additional-options -fdump-tree-gimple } */ + +#include "pr119692-1-1.C" + +/* { dg-bogus {using 'vector_length \(32\)', ignoring 1} {} { target openacc_nvidia_accel_selected xfail *-*-* } 0 } */ + +/* { dg-final { scan-tree-dump-not {(?n)#pragma omp target oacc_serial default\(none\) map\(tofrom:_ZTI2C2 \[len: [0-9]+\]\) map\(tofrom:_ZTI2C1 \[len: [0-9]+\]\) map\(tofrom:_ZTV2C1 \[len: [0-9]+\]\)$} gimple { xfail *-*-* } } } */ diff --git a/libgomp/testsuite/libgomp.oacc-c++/pr119692-1-3.C b/libgomp/testsuite/libgomp.oacc-c++/pr119692-1-3.C new file mode 100644 index 0000000..e9b44de --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/pr119692-1-3.C @@ -0,0 +1,12 @@ +/* PR119692 "C++ 'typeinfo', 'vtable' vs. OpenACC, OpenMP 'target' offloading" */ + +/* { dg-additional-options -DDEFAULT=default(present) } + Wrong code for offloading execution. + { dg-xfail-run-if PR119692 { ! openacc_host_selected } } */ +/* { dg-additional-options -fdump-tree-gimple } */ + +#include "pr119692-1-1.C" + +/* { dg-bogus {using 'vector_length \(32\)', ignoring 1} {} { target openacc_nvidia_accel_selected xfail *-*-* } 0 } */ + +/* { dg-final { scan-tree-dump-not {(?n)#pragma omp target oacc_serial default\(present\) map\(force_present:_ZTI2C2 \[len: [0-9]+\]\) map\(force_present:_ZTI2C1 \[len: [0-9]+\]\) map\(force_present:_ZTV2C1 \[len: [0-9]+\]\)$} gimple { xfail *-*-* } } } */ |