Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Replace SFINAE by if constexpr for create_mirror* #6

Open
wants to merge 136 commits into
base: refactor/create-mirror/variables
Choose a base branch
from
Open
Changes from 4 commits
Commits
Show all changes
136 commits
Select commit Hold shift + click to select a range
b4bc406
Reenable TestHIP_Memory_Requirements
Rombur Feb 2, 2024
c1a8006
Don't use Fedora development version in GitHub CI
masterleinad Mar 13, 2024
841b3a9
Fix deep copy when filling Rank-7 views
cedricchevalier19 Feb 15, 2024
a2f2ba4
TestViewCopy_c.hpp: add new unit test for deep copy (ViewFill)
cedricchevalier19 Feb 16, 2024
ae4d001
TestViewCopy_c.hpp: better handling for OpenMPTarget
cedricchevalier19 Mar 11, 2024
46354d2
Use builtin for atomic_fetch in the HIP backend
Rombur Mar 18, 2024
70603df
Combine the two Impl::create_mirror functions into one with constexpr
thierryantoun Mar 21, 2024
6524e2b
Format create_mirror
thierryantoun Mar 21, 2024
1635ff8
Combine the four Impl::create_mirror_view in one using if constexpr
thierryantoun Mar 21, 2024
d3bb03a
Combine the two create_mirror_view_and_copy functions into one with c…
pzehner Mar 21, 2024
872dc42
Fix Makefile.kokkos for Threads
masterleinad Mar 22, 2024
2035e31
Fix a bug in Makefile when using AMD GPU architectures (#6892)
Rombur Mar 25, 2024
8d734b0
Cuda: Fix configuring with CMake 3.28.4 (#6898)
masterleinad Mar 28, 2024
a53d30a
Merge pull request #6896 from masterleinad/fix_makefile_threads
dalg24 Mar 28, 2024
68c6684
Update Intel GPU architectures in Makefile (#6895)
masterleinad Mar 28, 2024
e2cfdec
Drop Experimental::LayoutTiled class template
dalg24 Mar 29, 2024
51b98e1
Get rid of now unnecessary use of is_layouttiled trait
dalg24 Mar 29, 2024
1efeb5d
Deprecate is_layouttiled trait
dalg24 Mar 29, 2024
5eac0bc
Merge pull request #6876 from masterleinad/disable_fedora_rawhide
dalg24 Apr 1, 2024
6355510
Move `Kokkos::Array` tests to a more suitable place (#6905)
dalg24 Apr 1, 2024
391e040
Do not return a copy of the input functor for Kokkos::Experimental::f…
tpadioleau Apr 2, 2024
b667853
Drop specialization of ViewMapping for Kokkos::Array
dalg24 Mar 29, 2024
059cd15
Accommodate users that depend on a code that define silly macros (#6909)
dalg24 Apr 2, 2024
2aecb1d
SYCL: Fix multi-GPU support and add test (#6887)
masterleinad Apr 3, 2024
5cf0951
Merge pull request #6910 from tpadioleau/remove-return-functor-copy-f…
dalg24 Apr 3, 2024
caa139c
SYCL: Unroll shuffle loops for top-level parallel_reduce and parallel…
masterleinad Apr 3, 2024
9220154
Merge branch 'refactor/create-mirror/variables' into refactor/create-…
pzehner Apr 4, 2024
a833fb0
Preparing readme for develop as the default branch (#6796)
cedricchevalier19 Apr 4, 2024
682291f
Format with clang
thierryantoun Apr 4, 2024
b89d618
Merge branch 'refactor/create-mirror/constexpr' of https://github.com…
thierryantoun Apr 4, 2024
497b438
CHANGELOG.md: 4.3.00 update
ndellingwood Apr 4, 2024
cc21a54
Merge pull request #6919 from ndellingwood/dev-changelog-4300
dalg24 Apr 4, 2024
4b90930
Refactor: Uniformize `create_mirror*` parameter name for views (#6917)
pzehner Apr 5, 2024
1256f69
Merge pull request #6822 from CExA-project/fix-deep-copy
dalg24 Apr 5, 2024
98b1a38
SYCL: Improve team_reduce implementation (#6562)
masterleinad Apr 5, 2024
e93b168
Merge pull request #6907 from dalg24/rm_experimental_layout_tiled
dalg24 Apr 6, 2024
e52cda3
Merge pull request #6785 from Rombur/memory_test
dalg24 Apr 8, 2024
55c5757
Use recommended/max team size functions in Cuda ParallelFor and Reduc…
tcclevenger Apr 8, 2024
8cf8410
SYCL: Fix range in subgroup scan for workgroup_scan
masterleinad Apr 8, 2024
7b41536
Merge pull request #6924 from masterleinad/fix_sycl_workgroup_scan
crtrott Apr 9, 2024
3a27cdb
Add ROCm 6.0 in the nightly CI
Rombur Apr 10, 2024
1fe8108
Merge pull request #6906 from dalg24/make_view_of_arrays_less_special
dalg24 Apr 10, 2024
74c8122
Merge pull request #6926 from Rombur/latest_rocm
dalg24 Apr 10, 2024
164519d
MI300 support unified memory support (#6877)
crtrott Apr 10, 2024
6ea7be7
cuda: reduction with `RangePolicy`: fix grid dimensions to work for l…
fnrizzi Apr 10, 2024
0099c10
Fix nightly CI
Rombur Apr 11, 2024
b0c2566
Merge pull request #6930 from Rombur/fix_nightly
dalg24 Apr 11, 2024
a2af4e0
Deprecate trailing Proxy template argument in Kokkos::Array
dalg24 Apr 11, 2024
92e02b5
CUDA: Update nvcc_wrapper
Apr 12, 2024
b5ec79b
Merge pull request #6936 from rgayatri23/issue_6874
dalg24 Apr 12, 2024
d88e2a5
bring back --fmad option to nvcc_wrapper (#6931)
glesur Apr 12, 2024
de3a263
Merge pull request #6934 from dalg24/deprecate_kokkos_array_proxy_tem…
dalg24 Apr 15, 2024
f2d3780
Remove unnecessary header include
dalg24 Apr 15, 2024
8c7cc95
Merge pull request #6940 from dalg24/unused_limits_header_include_in_…
dalg24 Apr 16, 2024
a8115e5
Adding converting constructor in Kokkos::RandomAccessIterator (#6929)
yasahi-hpc Apr 16, 2024
0e3a673
Use if constexpr for offset view create_mirror*
pzehner Apr 16, 2024
856bcc2
Use if constexpr for dynamic view create_mirror*
pzehner Apr 16, 2024
f1f4741
Use if constexpr for dynamic rank view create_mirror*
pzehner Apr 16, 2024
f94e8d3
Prefer standard C++ feature testing to guard the C++20 requires expre…
dalg24 Apr 16, 2024
c9e21ce
Add `kokkos_swap(Array<T, N>)` sepcialization
dalg24 Apr 16, 2024
730d8d8
Deprecate specialization of Kokkos::pair for a single element
dalg24 Apr 17, 2024
906e8ce
Merge pull request #6942 from dalg24/fix_nightlies_cxx20_requires_exp…
dalg24 Apr 17, 2024
d914fe3
Fix deprecated warning from `Kokkos::Array` specialization (#6945)
dalg24 Apr 17, 2024
e1b8afa
Add comments
pzehner Apr 17, 2024
1287f7f
Restore inline specifiers
pzehner Apr 17, 2024
69c527a
[ci skip] Enable deprecated code and deprecated warnings in nightly CI
Rombur Apr 17, 2024
e7b486f
Serial: Use the provided execution space instance in TeamPolicy
masterleinad Apr 17, 2024
0859ab0
Fixed the link for P6601 (Threads backend change)
nliber Apr 17, 2024
d5fd512
Merge pull request #6947 from dalg24/deprecate_kokkos_pair_void_speci…
dalg24 Apr 17, 2024
04bc3d9
Merge pull request #6952 from nliber/changelog43
crtrott Apr 18, 2024
34d0db2
Add test
masterleinad Apr 18, 2024
44fde21
Use Kokkos::AUTO for OpenMPTarget
masterleinad Apr 18, 2024
8706b68
kokkos_swap(Array) member friend should not be templated on some othe…
dalg24 Apr 18, 2024
86f5988
Fix noexcept specification for kokkos_swap on zero-sized arrays
dalg24 Apr 18, 2024
cc60295
Merge pull request #6951 from masterleinad/fix_serial_space_team_policy
crtrott Apr 19, 2024
3dc6f55
Add maybe_unused
pzehner Apr 22, 2024
2b54e2c
Mutualize check functions
pzehner Apr 22, 2024
7b6edf7
Simplify code
pzehner Apr 22, 2024
1ee800a
Fix missing maybe_unused
pzehner Apr 22, 2024
205fd15
Replace deprecated sycl::device_ptr/sycl::host_ptr
masterleinad Apr 22, 2024
5932685
Introduce alias based on feature macro
masterleinad Apr 22, 2024
a782773
Kokkos::Impl::SYCLTypes:: -> Kokkos::Impl::sycl_
masterleinad Apr 22, 2024
e2b7bb9
Merge pull request #6958 from masterleinad/sycl_replace_deprecated_us…
crtrott Apr 23, 2024
cf59f31
Merge pull request #6943 from dalg24/kokkos_swap_specialization_for_k…
crtrott Apr 23, 2024
ab3cae4
Fix wrong macro guards for deprecated Kokkos::pair<T1,void> specializ…
dalg24 Apr 24, 2024
fafe861
Fix support for Kokkos::Array of const-qualified element type
dalg24 Apr 24, 2024
2e82fdd
Merge pull request #6961 from dalg24/fixup_deprcated_guards_pair_void
dalg24 Apr 24, 2024
63eef46
Try to fix the CUDA 11.0 build
dalg24 Apr 24, 2024
ebb1cb3
Revert "Try to fix the CUDA 11.0 build"
dalg24 Apr 25, 2024
031f6d9
Alternate definition of Impl::is_nothrow_swappable_v for NVCC version…
dalg24 Apr 25, 2024
2391f17
Avoid introducing a 2nd definition of the Impl::swappable trait
dalg24 Apr 25, 2024
d434f87
Do not require OpenMP support for languages other than CXX
dalg24 Apr 25, 2024
19ca9ce
Update version
crtrott Apr 26, 2024
9686392
Add Linux Foundation notice and fix C++ standard
crtrott Apr 26, 2024
1864287
Merge pull request #6967 from crtrott/update-readme-kk-version
crtrott Apr 26, 2024
7e7709f
SYCL: Avoid deprecated floating-point number abs overloads (#6959)
masterleinad Apr 27, 2024
4ec8296
OpenMPTarget: Update loop order in MDRange (#6925)
rgayatri23 Apr 28, 2024
4fd1864
Restore previous types when create_mirror_view returns the source view
pzehner Apr 29, 2024
0140760
Fix linting
pzehner Apr 29, 2024
5ad1ff4
Remove unused namespaces
pzehner Apr 29, 2024
775be75
Merge branch 'develop' into refactor/create-mirror/constexpr
pzehner Apr 29, 2024
9eb9507
Merge branch 'develop' into refactor/create-mirror/constexpr
pzehner Apr 30, 2024
77ea52f
Threads: Don't silently allow m_instance to be a nullptr (#6969)
masterleinad May 1, 2024
4f416f3
Merge pull request #6965 from dalg24/cmake_openmp_cxx
dalg24 May 1, 2024
f699a2c
Fix enabling OpenMP with HIP and "compile as CMake language"
dalg24 May 1, 2024
2574b80
Fix OpenMP+CUDA when `Kokkos_ENABLE_COMPILE_AS_CMAKE_LANGUAGE` is `ON`
dalg24 May 1, 2024
27b3ced
Merge pull request #6949 from Rombur/nightly_deprecated
dalg24 May 1, 2024
15d13f2
Merge pull request #6882 from Rombur/hip_atomic_fetch
dalg24 May 1, 2024
ed4d254
Merge pull request #6972 from dalg24/fix_kokkos_compile_language_cuda…
crtrott May 1, 2024
dbd7f58
Merge pull request #6962 from dalg24/kokkos_array_const_qualified_ele…
crtrott May 1, 2024
1c8b624
Merge branch 'develop' into refactor/create-mirror/constexpr
pzehner May 2, 2024
ccd0126
Fix fedora CI builds with flang-new
masterleinad May 1, 2024
45a1404
Fix Copyright file
crtrott May 2, 2024
85610f4
Merge pull request #6984 from crtrott/Copyright
crtrott May 2, 2024
a75dc70
Merge pull request #6982 from masterleinad/fix_fedora
crtrott May 2, 2024
c6d8647
Also use is_nothrow_swappable workaround for Intel Classic Compilers …
masterleinad May 2, 2024
b29b14f
Merge branch 'develop' into refactor/create-mirror/constexpr
pzehner May 3, 2024
69567f3
Add thread-safety tests (#6938)
masterleinad May 3, 2024
9c79202
Fix deprecation warnings with GCC for pair<T1,void> comparison operators
dalg24 May 3, 2024
7b8e3a6
Fix TPL_LIBRARY_SUFFIXES for 32-bit build
masterleinad May 6, 2024
2826017
Avoid duplicated definition of KOKKOS_IMPL_32BIT
masterleinad May 6, 2024
ccadc7d
Disable failing parallel_scan_with_reducers test
masterleinad May 6, 2024
06e4c5b
Merge pull request #6989 from dalg24/deprecated_attribute_comparison_…
dalg24 May 6, 2024
e4cc686
Merge pull request #6990 from masterleinad/fix_32bit_tpl_library_path
dalg24 May 7, 2024
d61d75a
Fix a bug when using realloc on views of non-default constructible el…
aprokop May 8, 2024
50a862c
SYCL: Prepare Parallel* for Graphs (#6988)
masterleinad May 8, 2024
f5b3422
SYCL: Fix deprecation in custom parallel_for RangePolicy implementation
masterleinad May 8, 2024
37986fd
[ci skip] update changelog for 4.3.1 (#6995)
ndellingwood May 8, 2024
7cad3e7
OpenMPTarget: Use mutex lock for parallel scan.
May 8, 2024
a69e81a
Merge pull request #6998 from rgayatri23/ompt_scan_lock
dalg24 May 9, 2024
5a5306c
Merge pull request #6997 from masterleinad/sycl_fix_custom_parallel_f…
dalg24 May 9, 2024
00170ae
Remove cuSPARSE TPL
dalg24 May 9, 2024
506da18
Merge pull request #7002 from dalg24/rm_tpl_cusparse
dalg24 May 9, 2024
1d9d0df
SYCL: Print submission command queue property (#7004)
masterleinad May 10, 2024
df018d9
Suppress deprecated warnings via pragma push/pop in the tests (#6999)
dalg24 May 13, 2024
be07ebc
Merge branch 'develop' into refactor/create-mirror/constexpr
pzehner May 14, 2024
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
204 changes: 72 additions & 132 deletions core/src/Kokkos_CopyViews.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -3471,35 +3471,22 @@ void check_view_ctor_args_create_mirror() {
}

template <class T, class... P, class... ViewCtorArgs>
inline std::enable_if_t<!Impl::ViewCtorProp<ViewCtorArgs...>::has_memory_space,
typename Kokkos::View<T, P...>::HostMirror>
create_mirror(const Kokkos::View<T, P...>& src,
const Impl::ViewCtorProp<ViewCtorArgs...>& arg_prop) {
using src_type = View<T, P...>;
using dst_type = typename src_type::HostMirror;

check_view_ctor_args_create_mirror<ViewCtorArgs...>();

auto prop_copy = Impl::with_properties_if_unset(
arg_prop, std::string(src.label()).append("_mirror"));

return dst_type(prop_copy, src.layout());
}

// Create a mirror in a new space (specialization for different space)
template <class T, class... P, class... ViewCtorArgs,
class Enable = std::enable_if_t<
Impl::ViewCtorProp<ViewCtorArgs...>::has_memory_space>>
auto create_mirror(const Kokkos::View<T, P...>& src,
const Impl::ViewCtorProp<ViewCtorArgs...>& arg_prop) {
check_view_ctor_args_create_mirror<ViewCtorArgs...>();

auto prop_copy = Impl::with_properties_if_unset(
arg_prop, std::string(src.label()).append("_mirror"));
using alloc_prop = decltype(prop_copy);

return typename Impl::MirrorType<typename alloc_prop::memory_space, T,
P...>::view_type(prop_copy, src.layout());
if constexpr (Impl::ViewCtorProp<ViewCtorArgs...>::has_memory_space) {
using memory_space = typename decltype(prop_copy)::memory_space;
using dst_type =
typename Impl::MirrorType<memory_space, T, P...>::view_type;
return dst_type(prop_copy, src.layout());
} else {
using dst_type = typename View<T, P...>::HostMirror;
return dst_type(prop_copy, src.layout());
}
}
} // namespace Impl

Expand Down Expand Up @@ -3551,71 +3538,37 @@ std::enable_if_t<std::is_void<typename ViewTraits<T, P...>::specialize>::value,
typename Impl::MirrorType<Space, T, P...>::view_type>
create_mirror(Kokkos::Impl::WithoutInitializing_t wi, Space const&,
Kokkos::View<T, P...> const& src) {
return Impl::create_mirror(src, view_alloc(typename Space::memory_space{}, wi));
return Impl::create_mirror(src,
view_alloc(typename Space::memory_space{}, wi));
tpadioleau marked this conversation as resolved.
Show resolved Hide resolved
}

namespace Impl {

template <class T, class... P, class... ViewCtorArgs>
inline std::enable_if_t<
!Impl::ViewCtorProp<ViewCtorArgs...>::has_memory_space &&
(std::is_same<
typename Kokkos::View<T, P...>::memory_space,
typename Kokkos::View<T, P...>::HostMirror::memory_space>::value &&
std::is_same<
typename Kokkos::View<T, P...>::data_type,
typename Kokkos::View<T, P...>::HostMirror::data_type>::value),
typename Kokkos::View<T, P...>::HostMirror>
create_mirror_view(const Kokkos::View<T, P...>& src,
const Impl::ViewCtorProp<ViewCtorArgs...>&) {
check_view_ctor_args_create_mirror<ViewCtorArgs...>();
return src;
}

template <class T, class... P, class... ViewCtorArgs>
inline std::enable_if_t<
!Impl::ViewCtorProp<ViewCtorArgs...>::has_memory_space &&
!(std::is_same<typename Kokkos::View<T, P...>::memory_space,
typename Kokkos::View<
T, P...>::HostMirror::memory_space>::value &&
std::is_same<
typename Kokkos::View<T, P...>::data_type,
typename Kokkos::View<T, P...>::HostMirror::data_type>::value),
typename Kokkos::View<T, P...>::HostMirror>
create_mirror_view(const Kokkos::View<T, P...>& src,
const Impl::ViewCtorProp<ViewCtorArgs...>& arg_prop) {
return Kokkos::Impl::create_mirror(src, arg_prop);
}

// Create a mirror view in a new space (specialization for same space)
template <class T, class... P, class... ViewCtorArgs,
class = std::enable_if_t<
Impl::ViewCtorProp<ViewCtorArgs...>::has_memory_space>>
std::enable_if_t<Impl::MirrorViewType<
typename Impl::ViewCtorProp<ViewCtorArgs...>::memory_space,
T, P...>::is_same_memspace,
typename Impl::MirrorViewType<
typename Impl::ViewCtorProp<ViewCtorArgs...>::memory_space,
T, P...>::view_type>
create_mirror_view(const Kokkos::View<T, P...>& src,
const Impl::ViewCtorProp<ViewCtorArgs...>&) {
check_view_ctor_args_create_mirror<ViewCtorArgs...>();
return src;
}

// Create a mirror view in a new space (specialization for different space)
template <class T, class... P, class... ViewCtorArgs,
class = std::enable_if_t<
Impl::ViewCtorProp<ViewCtorArgs...>::has_memory_space>>
std::enable_if_t<!Impl::MirrorViewType<
typename Impl::ViewCtorProp<ViewCtorArgs...>::memory_space,
T, P...>::is_same_memspace,
typename Impl::MirrorViewType<
typename Impl::ViewCtorProp<ViewCtorArgs...>::memory_space,
T, P...>::view_type>
create_mirror_view(const Kokkos::View<T, P...>& src,
const Impl::ViewCtorProp<ViewCtorArgs...>& arg_prop) {
return Kokkos::Impl::create_mirror(src, arg_prop);
auto create_mirror_view(const Kokkos::View<T, P...>& src,
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I notice two changes:

  • the use of auto,
  • some functions were marked inline.

I would suggest to mention these changes when you open a PR.

Copy link
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The inline functions were also templated, so I wonder if the keyword was necessary.

Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Well the keyword is primarily a hint for the compiler to actually inline a function call before avoiding an ODR violation.

Personally I would remove the keyword in this case.

Copy link
Member Author

@pzehner pzehner Apr 19, 2024

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I changed my mind about the inline keyword, and let its use unaltered. According to this SO answer, function templates are not inlined by default. This change was outside the scope of this PR.

const Impl::ViewCtorProp<ViewCtorArgs...>& arg_prop) {
if constexpr (!Impl::ViewCtorProp<ViewCtorArgs...>::has_memory_space) {
if constexpr (std::is_same<typename Kokkos::View<T, P...>::memory_space,
typename Kokkos::View<
T, P...>::HostMirror::memory_space>::value &&
tpadioleau marked this conversation as resolved.
Show resolved Hide resolved
std::is_same<typename Kokkos::View<T, P...>::data_type,
typename Kokkos::View<
T, P...>::HostMirror::data_type>::value) {
check_view_ctor_args_create_mirror<ViewCtorArgs...>();
return src;
} else {
return Kokkos::Impl::create_mirror(src, arg_prop);
}
} else {
if constexpr (Impl::MirrorViewType<typename Impl::ViewCtorProp<
ViewCtorArgs...>::memory_space,
T, P...>::is_same_memspace) {
check_view_ctor_args_create_mirror<ViewCtorArgs...>();
return src;
} else {
return Kokkos::Impl::create_mirror(src, arg_prop);
}
}
}
} // namespace Impl

Expand Down Expand Up @@ -3686,16 +3639,12 @@ auto create_mirror_view(const Impl::ViewCtorProp<ViewCtorArgs...>& arg_prop,
return Impl::create_mirror_view(src, arg_prop);
}

template <class... ViewCtorArgs, class T, class... P>
auto create_mirror_view_and_copy(
const Impl::ViewCtorProp<ViewCtorArgs...>&,
const Kokkos::View<T, P...>& src,
std::enable_if_t<
std::is_void<typename ViewTraits<T, P...>::specialize>::value &&
Impl::MirrorViewType<
typename Impl::ViewCtorProp<ViewCtorArgs...>::memory_space, T,
P...>::is_same_memspace>* = nullptr) {
namespace Impl {

template <class... ViewCtorArgs>
void check_view_ctor_args_create_mirror_view_and_copy() {
using alloc_prop_input = Impl::ViewCtorProp<ViewCtorArgs...>;

static_assert(
alloc_prop_input::has_memory_space,
"The view constructor arguments passed to "
Expand All @@ -3708,52 +3657,43 @@ auto create_mirror_view_and_copy(
"The view constructor arguments passed to "
"Kokkos::create_mirror_view_and_copy must "
"not explicitly allow padding!");

// same behavior as deep_copy(src, src)
if (!alloc_prop_input::has_execution_space)
fence(
"Kokkos::create_mirror_view_and_copy: fence before returning src view");
return src;
}

template <class... ViewCtorArgs, class T, class... P>
} // namespace Impl

template <class... ViewCtorArgs, class T, class... P,
class = std::enable_if<std::is_void<typename ViewTraits<T, P...>::specialize>::value>>
auto create_mirror_view_and_copy(
const Impl::ViewCtorProp<ViewCtorArgs...>& arg_prop,
const Kokkos::View<T, P...>& src,
std::enable_if_t<
std::is_void<typename ViewTraits<T, P...>::specialize>::value &&
!Impl::MirrorViewType<
typename Impl::ViewCtorProp<ViewCtorArgs...>::memory_space, T,
P...>::is_same_memspace>* = nullptr) {
const Kokkos::View<T, P...>& src) {
using alloc_prop_input = Impl::ViewCtorProp<ViewCtorArgs...>;
static_assert(
alloc_prop_input::has_memory_space,
"The view constructor arguments passed to "
"Kokkos::create_mirror_view_and_copy must include a memory space!");
static_assert(!alloc_prop_input::has_pointer,
"The view constructor arguments passed to "
"Kokkos::create_mirror_view_and_copy must "
"not include a pointer!");
static_assert(!alloc_prop_input::allow_padding,
"The view constructor arguments passed to "
"Kokkos::create_mirror_view_and_copy must "
"not explicitly allow padding!");
using Space = typename alloc_prop_input::memory_space;
using Mirror = typename Impl::MirrorViewType<Space, T, P...>::view_type;

auto arg_prop_copy = Impl::with_properties_if_unset(
arg_prop, std::string{}, WithoutInitializing,
typename Space::execution_space{});

std::string& label = Impl::get_property<Impl::LabelTag>(arg_prop_copy);
if (label.empty()) label = src.label();
auto mirror = typename Mirror::non_const_type{arg_prop_copy, src.layout()};
if constexpr (alloc_prop_input::has_execution_space) {
deep_copy(Impl::get_property<Impl::ExecutionSpaceTag>(arg_prop_copy),
mirror, src);
} else
deep_copy(mirror, src);
return mirror;

Impl::check_view_ctor_args_create_mirror_view_and_copy<ViewCtorArgs...>();

if constexpr (Impl::MirrorViewType<typename alloc_prop_input::memory_space, T, P...>::is_same_memspace) {
// same behavior as deep_copy(src, src)
if constexpr (!alloc_prop_input::has_execution_space)
fence(
"Kokkos::create_mirror_view_and_copy: fence before returning src view");
return src;
} else {
using Space = typename alloc_prop_input::memory_space;
using Mirror = typename Impl::MirrorViewType<Space, T, P...>::view_type;

auto arg_prop_copy = Impl::with_properties_if_unset(
arg_prop, std::string{}, WithoutInitializing,
typename Space::execution_space{});

std::string& label = Impl::get_property<Impl::LabelTag>(arg_prop_copy);
if (label.empty()) label = src.label();
auto mirror = typename Mirror::non_const_type{arg_prop_copy, src.layout()};
if constexpr (alloc_prop_input::has_execution_space) {
deep_copy(Impl::get_property<Impl::ExecutionSpaceTag>(arg_prop_copy),
mirror, src);
} else
deep_copy(mirror, src);
return mirror;
}
}

// Previously when using auto here, the intel compiler 19.3 would
Expand Down