1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
|
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
// SPDX-FileCopyrightText: Copyright Contributors to the Kokkos project
#ifndef KOKKOS_IMPL_PUBLIC_INCLUDE
#define KOKKOS_IMPL_PUBLIC_INCLUDE
#endif
#include <HIP/Kokkos_HIP_DeepCopy.hpp>
#include <HIP/Kokkos_HIP_Error.hpp> // HIP_SAFE_CALL
#include <HIP/Kokkos_HIP.hpp>
#include <HIP/Kokkos_HIP_Instance.hpp>
namespace Kokkos {
namespace Impl {
namespace {
hipStream_t get_deep_copy_stream() {
static hipStream_t s = nullptr;
if (s == nullptr) {
KOKKOS_IMPL_HIP_SAFE_CALL(hipStreamCreate(&s));
}
return s;
}
} // namespace
void DeepCopyHIP(void* dst, void const* src, size_t n) {
KOKKOS_IMPL_HIP_SAFE_CALL(hipMemcpyAsync(dst, src, n, hipMemcpyDefault));
}
void DeepCopyAsyncHIP(const HIP& instance, void* dst, void const* src,
size_t n) {
KOKKOS_IMPL_HIP_SAFE_CALL(
instance.impl_internal_space_instance()->hip_memcpy_async_wrapper(
dst, src, n, hipMemcpyDefault));
}
void DeepCopyAsyncHIP(void* dst, void const* src, size_t n) {
hipStream_t s = get_deep_copy_stream();
KOKKOS_IMPL_HIP_SAFE_CALL(hipMemcpyAsync(dst, src, n, hipMemcpyDefault, s));
Kokkos::Tools::Experimental::Impl::profile_fence_event<HIP>(
"Kokkos::Impl::DeepCopyAsyncHIP: Post Deep Copy Fence on Deep-Copy "
"stream",
Kokkos::Tools::Experimental::SpecialSynchronizationCases::
DeepCopyResourceSynchronization,
[&]() { KOKKOS_IMPL_HIP_SAFE_CALL(hipStreamSynchronize(s)); });
}
} // namespace Impl
} // namespace Kokkos
|
