Fix late-coming issues. (#1724)
Implementation for hipMemcpyWithStream.
This commit is contained in:
Alex Voicu
committed by
Maneesh Gupta
Maneesh Gupta
parent
32625638da
commit
150e690a3a
@@ -48,8 +48,6 @@ typedef struct grid_launch_parm
|
||||
//! Value of packet fences to apply to launch.
|
||||
//! The correspond to the value of bits 9:14 in the AQL packet,
|
||||
//! see HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE and hsa_fence_scope_t.
|
||||
//! Set to -1 for conservative defaults.
|
||||
//! Placeholder, is not used to control packet dispatch yet
|
||||
unsigned int launch_fence;
|
||||
|
||||
//! Pointer to the accelerator_view where the kernel should execute.
|
||||
|
||||
@@ -1361,6 +1361,9 @@ hipError_t hipHostFree(void* ptr);
|
||||
*/
|
||||
hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind);
|
||||
|
||||
// TODO: Add description
|
||||
hipError_t hipMemcpyWithStream(void* dst, const void* src, size_t sizeBytes,
|
||||
hipMemcpyKind kind, hipStream_t stream);
|
||||
/**
|
||||
* @brief Copy data from Host to Device
|
||||
*
|
||||
|
||||
@@ -611,6 +611,18 @@ inline static hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes,
|
||||
}
|
||||
|
||||
|
||||
inline hipError_t hipMemcpyWithStream(void* dst, const void* src,
|
||||
size_t sizeBytes, hipMemcpyKind copyKind,
|
||||
hipStream_t stream) {
|
||||
cudaError_t error = cudaMemcpyAsync(dst, src, sizeBytes,
|
||||
hipMemcpyKindToCudaMemcpyKind(copyKind),
|
||||
stream);
|
||||
|
||||
if (error != cudaSuccess) return hipCUDAErrorTohipError(error);
|
||||
|
||||
return hipCUDAErrorTohipError(cudaStreamSynchronize(stream));
|
||||
}
|
||||
|
||||
inline static hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes,
|
||||
hipMemcpyKind copyKind, hipStream_t stream __dparm(0)) {
|
||||
return hipCUDAErrorTohipError(
|
||||
|
||||
+10
-8
@@ -1619,16 +1619,18 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, grid_
|
||||
lp->group_dim.y = block.y;
|
||||
lp->group_dim.z = block.z;
|
||||
lp->barrier_bit = barrier_bit_queue_default;
|
||||
lp->launch_fence = -1;
|
||||
|
||||
if (!lockAcquired) {
|
||||
auto crit = stream->lockopen_preKernelCommand();
|
||||
lp->av = &(crit->_av);
|
||||
} else {
|
||||
// this stream is already locked (e.g., call from hipExtLaunchMultiKernelMultiDevice)
|
||||
lp->av = &(stream->criticalData()._av);
|
||||
}
|
||||
if (!lockAcquired) stream->lockopen_preKernelCommand();
|
||||
auto &crit = stream->criticalData();
|
||||
lp->av = &(crit._av);
|
||||
lp->cf = nullptr;
|
||||
auto acq = (HCC_OPT_FLUSH && !crit._last_op_was_a_copy) ?
|
||||
HSA_FENCE_SCOPE_AGENT : HSA_FENCE_SCOPE_SYSTEM;
|
||||
auto rel = HCC_OPT_FLUSH ?
|
||||
HSA_FENCE_SCOPE_AGENT : HSA_FENCE_SCOPE_SYSTEM;
|
||||
lp->launch_fence = (acq << HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE) |
|
||||
(rel << HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE);
|
||||
crit._last_op_was_a_copy = false;
|
||||
ihipPrintKernelLaunch(kernelNameStr, lp, stream);
|
||||
|
||||
return (stream);
|
||||
|
||||
@@ -495,9 +495,10 @@ struct LockedBase {
|
||||
|
||||
template <typename MUTEX_TYPE>
|
||||
class ihipStreamCriticalBase_t : public LockedBase<MUTEX_TYPE> {
|
||||
public:
|
||||
public:
|
||||
ihipStreamCriticalBase_t(ihipStream_t* parentStream, hc::accelerator_view av)
|
||||
: _av(av), _parent(parentStream){};
|
||||
: _parent{parentStream}, _av{av}, _last_op_was_a_copy{false}
|
||||
{}
|
||||
|
||||
~ihipStreamCriticalBase_t() {}
|
||||
|
||||
@@ -519,12 +520,9 @@ class ihipStreamCriticalBase_t : public LockedBase<MUTEX_TYPE> {
|
||||
return gotLock ? this : nullptr;
|
||||
};
|
||||
|
||||
public:
|
||||
ihipStream_t* _parent;
|
||||
|
||||
hc::accelerator_view _av;
|
||||
|
||||
private:
|
||||
bool _last_op_was_a_copy;
|
||||
};
|
||||
|
||||
|
||||
|
||||
+386
-131
@@ -27,6 +27,7 @@ THE SOFTWARE.
|
||||
#include "hip_hcc_internal.h"
|
||||
#include "trace_helper.h"
|
||||
|
||||
#include <functional>
|
||||
#include <fstream>
|
||||
|
||||
__device__ char __hip_device_heap[__HIP_SIZE_OF_HEAP];
|
||||
@@ -35,23 +36,372 @@ __device__ uint32_t __hip_device_page_flag[__HIP_NUM_PAGES];
|
||||
// Internal HIP APIS:
|
||||
namespace hip_internal {
|
||||
|
||||
hipError_t memcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind,
|
||||
hipStream_t stream) {
|
||||
hipError_t e = hipSuccess;
|
||||
namespace {
|
||||
inline
|
||||
const char* hsa_to_string(hsa_status_t err) noexcept
|
||||
{
|
||||
const char* r{};
|
||||
|
||||
// Return success if number of bytes to copy is 0
|
||||
if (sizeBytes == 0) return e;
|
||||
if (hsa_status_string(err, &r) == HSA_STATUS_SUCCESS) return r;
|
||||
|
||||
return "Unknown.";
|
||||
}
|
||||
|
||||
template<std::size_t m, std::size_t n>
|
||||
inline
|
||||
void throwing_result_check(hsa_status_t res, const char (&file)[m],
|
||||
const char (&function)[n], int line) {
|
||||
if (res == HSA_STATUS_SUCCESS) return;
|
||||
if (res == HSA_STATUS_INFO_BREAK) return;
|
||||
|
||||
throw std::runtime_error{"Failed in file " + (file +
|
||||
(", in function \"" + (function +
|
||||
("\", on line " + std::to_string(line))))) +
|
||||
", with error: " + hsa_to_string(res)};
|
||||
}
|
||||
|
||||
inline
|
||||
hsa_agent_t cpu_agent() {
|
||||
hsa_agent_t r{};
|
||||
throwing_result_check(hsa_iterate_agents([](hsa_agent_t x, void* pr) {
|
||||
hsa_device_type_t t{};
|
||||
hsa_agent_get_info(x, HSA_AGENT_INFO_DEVICE, &t);
|
||||
|
||||
if (t != HSA_DEVICE_TYPE_CPU) return HSA_STATUS_SUCCESS;
|
||||
|
||||
*static_cast<hsa_agent_t *>(pr) = x;
|
||||
|
||||
return HSA_STATUS_INFO_BREAK;
|
||||
}, &r), __FILE__, __func__, __LINE__);
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
inline
|
||||
hsa_device_type_t type(hsa_agent_t x)
|
||||
{
|
||||
hsa_device_type_t r{};
|
||||
throwing_result_check(hsa_agent_get_info(x, HSA_AGENT_INFO_DEVICE, &r),
|
||||
__FILE__, __func__, __LINE__);
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
const auto is_large_BAR{[](){
|
||||
std::unique_ptr<void, void (*)(void*)> hsa{
|
||||
(hsa_init() == HSA_STATUS_SUCCESS)
|
||||
? reinterpret_cast<void*>(UINT64_MAX) : nullptr,
|
||||
[](void* p) { if (p) hsa_shut_down(); }};
|
||||
|
||||
if (!hsa) return false;
|
||||
|
||||
bool r{true};
|
||||
|
||||
throwing_result_check(hsa_iterate_agents([](hsa_agent_t x, void* pr) {
|
||||
if (x.handle == cpu_agent().handle) return HSA_STATUS_SUCCESS;
|
||||
|
||||
throwing_result_check(
|
||||
hsa_agent_iterate_regions(x, [](hsa_region_t y, void* p) {
|
||||
hsa_region_segment_t seg{};
|
||||
throwing_result_check(
|
||||
hsa_region_get_info(y, HSA_REGION_INFO_SEGMENT, &seg),
|
||||
__FILE__, __func__, __LINE__);
|
||||
|
||||
if (seg != HSA_REGION_SEGMENT_GLOBAL) {
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
uint32_t flags{};
|
||||
throwing_result_check(hsa_region_get_info(
|
||||
y, HSA_REGION_INFO_GLOBAL_FLAGS, &flags),
|
||||
__FILE__, __func__, __LINE__);
|
||||
|
||||
if (flags & HSA_REGION_GLOBAL_FLAG_COARSE_GRAINED) {
|
||||
hsa_amd_memory_pool_access_t tmp{};
|
||||
throwing_result_check(
|
||||
hsa_amd_agent_memory_pool_get_info(
|
||||
cpu_agent(),
|
||||
hsa_amd_memory_pool_t{y.handle},
|
||||
HSA_AMD_AGENT_MEMORY_POOL_INFO_ACCESS,
|
||||
&tmp),
|
||||
__FILE__, __func__, __LINE__);
|
||||
|
||||
*static_cast<bool*>(p) &=
|
||||
tmp != HSA_AMD_MEMORY_POOL_ACCESS_NEVER_ALLOWED;
|
||||
}
|
||||
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}, pr), __FILE__, __func__, __LINE__);
|
||||
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}, &r), __FILE__, __func__, __LINE__);
|
||||
|
||||
return r;
|
||||
}()};
|
||||
|
||||
inline
|
||||
hsa_amd_pointer_info_t info(const void* p)
|
||||
{
|
||||
hsa_amd_pointer_info_t r{sizeof(hsa_amd_pointer_info_t)};
|
||||
throwing_result_check(
|
||||
hsa_amd_pointer_info(
|
||||
const_cast<void*>(p), &r, nullptr, nullptr, nullptr),
|
||||
__FILE__, __func__, __LINE__);
|
||||
|
||||
r.size = is_large_BAR || (type(r.agentOwner) == HSA_DEVICE_TYPE_CPU) ?
|
||||
UINT32_MAX : sizeof(hsa_amd_pointer_info_t);
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
constexpr size_t staging_sz{4 * 1024 * 1024}; // 2 Pages.
|
||||
|
||||
thread_local const std::unique_ptr<void, void (*)(void *)> staging_buffer{
|
||||
[]() {
|
||||
hsa_region_t r{};
|
||||
throwing_result_check(hsa_agent_iterate_regions(
|
||||
cpu_agent(), [](hsa_region_t x, void *p) {
|
||||
hsa_region_segment_t seg{};
|
||||
throwing_result_check(
|
||||
hsa_region_get_info(x, HSA_REGION_INFO_SEGMENT, &seg),
|
||||
__FILE__, __func__, __LINE__);
|
||||
|
||||
if (seg != HSA_REGION_SEGMENT_GLOBAL) return HSA_STATUS_SUCCESS;
|
||||
|
||||
uint32_t flags{};
|
||||
throwing_result_check(hsa_region_get_info(
|
||||
x, HSA_REGION_INFO_GLOBAL_FLAGS, &flags),
|
||||
__FILE__, __func__, __LINE__);
|
||||
|
||||
if (flags & HSA_REGION_GLOBAL_FLAG_COARSE_GRAINED) {
|
||||
*static_cast<hsa_region_t *>(p) = x;
|
||||
|
||||
return HSA_STATUS_INFO_BREAK;
|
||||
}
|
||||
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}, &r), __FILE__, __func__, __LINE__);
|
||||
|
||||
void *tp{};
|
||||
throwing_result_check(hsa_memory_allocate(r, staging_sz, &tp),
|
||||
__FILE__, __func__, __LINE__);
|
||||
|
||||
return tp;
|
||||
}(),
|
||||
[](void *ptr) { hsa_memory_free(ptr); }};
|
||||
|
||||
thread_local hsa_signal_t copy_signal{[]() {
|
||||
hsa_agent_t cpu{cpu_agent()};
|
||||
hsa_signal_t sgn{};
|
||||
throwing_result_check(hsa_signal_create(1, 1, &cpu, &sgn),
|
||||
__FILE__, __func__, __LINE__);
|
||||
|
||||
return sgn;
|
||||
}()};
|
||||
} // Unnamed namespace.
|
||||
|
||||
inline
|
||||
void do_copy(void* __restrict dst, const void* __restrict src, std::size_t n,
|
||||
hsa_agent_t da, hsa_agent_t sa) {
|
||||
hsa_signal_silent_store_relaxed(copy_signal, 1);
|
||||
throwing_result_check(
|
||||
hsa_amd_memory_async_copy(dst, da, src, sa, n, 0, nullptr, copy_signal),
|
||||
__FILE__, __func__, __LINE__);
|
||||
|
||||
while (hsa_signal_wait_relaxed(copy_signal, HSA_SIGNAL_CONDITION_EQ, 0,
|
||||
UINT64_MAX, HSA_WAIT_STATE_ACTIVE));
|
||||
}
|
||||
|
||||
inline
|
||||
void do_std_memcpy(
|
||||
void* __restrict dst, const void* __restrict src, std::size_t n) {
|
||||
std::memcpy(dst, src, n);
|
||||
|
||||
return std::atomic_thread_fence(std::memory_order_seq_cst);
|
||||
}
|
||||
|
||||
inline
|
||||
void d2h_copy(void* __restrict dst, const void* __restrict src, size_t n,
|
||||
hsa_amd_pointer_info_t si) {
|
||||
// TODO: characterise direct largeBAR reads from agent-allocated memory.
|
||||
// if (si.size == UINT32_MAX) {
|
||||
// return do_std_memcpy(dst, src, n);
|
||||
// }
|
||||
|
||||
const auto di{info(dst)};
|
||||
|
||||
if (di.type == HSA_EXT_POINTER_TYPE_LOCKED) {
|
||||
dst = static_cast<char*>(di.agentBaseAddress) +
|
||||
(static_cast<char*>(dst) -
|
||||
static_cast<char*>(di.hostBaseAddress));
|
||||
do_copy(dst, src, n, si.agentOwner, si.agentOwner);
|
||||
}
|
||||
else if (n <= staging_sz) {
|
||||
do_copy(staging_buffer.get(), src, n, si.agentOwner, si.agentOwner);
|
||||
std::memcpy(dst, staging_buffer.get(), n);
|
||||
}
|
||||
else {
|
||||
std::unique_ptr<void, void (*)(void*)> lck{
|
||||
dst, [](void* p) { hsa_amd_memory_unlock(p); }};
|
||||
|
||||
throwing_result_check(hsa_amd_memory_lock(dst, n, &si.agentOwner, 1,
|
||||
const_cast<void**>(&dst)),
|
||||
__FILE__, __func__, __LINE__);
|
||||
|
||||
do_copy(dst, src, n, si.agentOwner, si.agentOwner);
|
||||
}
|
||||
}
|
||||
|
||||
inline
|
||||
void h2d_copy(void* __restrict dst, const void* __restrict src, size_t n,
|
||||
hsa_amd_pointer_info_t di) {
|
||||
if (di.size == UINT32_MAX) {
|
||||
return do_std_memcpy(dst, src, n);
|
||||
}
|
||||
|
||||
const auto si{info(const_cast<void*>(src))};
|
||||
|
||||
if (si.type == HSA_EXT_POINTER_TYPE_LOCKED) {
|
||||
src = static_cast<char*>(si.agentBaseAddress) +
|
||||
(static_cast<const char*>(src) -
|
||||
static_cast<char*>(si.hostBaseAddress));
|
||||
do_copy(dst, src, n, di.agentOwner, di.agentOwner);
|
||||
}
|
||||
else if (n <= staging_sz) {
|
||||
std::memcpy(staging_buffer.get(), src, n);
|
||||
do_copy(dst, staging_buffer.get(), n, di.agentOwner, di.agentOwner);
|
||||
}
|
||||
else {
|
||||
std::unique_ptr<void, void (*)(void*)> lck{
|
||||
const_cast<void*>(src), [](void* p) { hsa_amd_memory_unlock(p); }};
|
||||
|
||||
throwing_result_check(hsa_amd_memory_lock(const_cast<void*>(src), n,
|
||||
&di.agentOwner, 1,
|
||||
const_cast<void**>(&src)),
|
||||
__FILE__, __func__, __LINE__);
|
||||
|
||||
do_copy(dst, src, n, di.agentOwner, di.agentOwner);
|
||||
}
|
||||
}
|
||||
|
||||
inline
|
||||
void generic_copy(void* __restrict dst, const void* __restrict src, size_t n,
|
||||
hsa_amd_pointer_info_t di, hsa_amd_pointer_info_t si) {
|
||||
if (di.size == UINT32_MAX && si.size == UINT32_MAX) {
|
||||
return do_std_memcpy(dst, src, n);
|
||||
}
|
||||
|
||||
std::unique_ptr<void, void (*)(void*)> lck0{
|
||||
nullptr, [](void* p) { hsa_amd_memory_unlock(p); }};
|
||||
std::unique_ptr<void, void (*)(void*)> lck1{nullptr, lck0.get_deleter()};
|
||||
|
||||
switch (si.type) {
|
||||
case HSA_EXT_POINTER_TYPE_HSA:
|
||||
if (di.type == HSA_EXT_POINTER_TYPE_HSA) {
|
||||
hsa_memory_copy(dst, src, n);
|
||||
return; // TODO: do_copy(dst, src, n, di.agentOwner, si.agentOwner);
|
||||
}
|
||||
|
||||
if (di.type == HSA_EXT_POINTER_TYPE_UNKNOWN ||
|
||||
di.type == HSA_EXT_POINTER_TYPE_LOCKED) {
|
||||
return d2h_copy(dst, src, n, si);
|
||||
}
|
||||
break;
|
||||
case HSA_EXT_POINTER_TYPE_LOCKED:
|
||||
if (di.type == HSA_EXT_POINTER_TYPE_UNKNOWN) {
|
||||
std::memcpy(dst, si.hostBaseAddress, n);
|
||||
|
||||
return;
|
||||
}
|
||||
if (di.type == HSA_EXT_POINTER_TYPE_LOCKED) {
|
||||
std::memcpy(di.hostBaseAddress, si.hostBaseAddress, n);
|
||||
|
||||
return;
|
||||
}
|
||||
src = si.agentBaseAddress;
|
||||
si.agentOwner = di.agentOwner;
|
||||
break;
|
||||
case HSA_EXT_POINTER_TYPE_UNKNOWN:
|
||||
if (di.type == HSA_EXT_POINTER_TYPE_UNKNOWN) {
|
||||
std::memcpy(dst, src, n);
|
||||
|
||||
return;
|
||||
}
|
||||
if (di.type == HSA_EXT_POINTER_TYPE_LOCKED) {
|
||||
std::memcpy(di.hostBaseAddress, src, n);
|
||||
|
||||
return;
|
||||
}
|
||||
return h2d_copy(dst, src, n, di);
|
||||
default: do_copy(dst, src, n, di.agentOwner, si.agentOwner); break;
|
||||
}
|
||||
}
|
||||
|
||||
inline
|
||||
void memcpy_impl(void* __restrict dst, const void* __restrict src, size_t n,
|
||||
hipMemcpyKind k) noexcept {
|
||||
switch (k) {
|
||||
case hipMemcpyHostToHost: std::memcpy(dst, src, n); break;
|
||||
case hipMemcpyHostToDevice:
|
||||
return is_large_BAR ? do_std_memcpy(dst, src, n)
|
||||
: h2d_copy(dst, src, n, info(dst));
|
||||
case hipMemcpyDeviceToHost:
|
||||
// TODO: characterise direct largeBAR reads from agent-allocated memory.
|
||||
return /*is_large_BAR ? do_std_memcpy(dst, src, n)
|
||||
: */d2h_copy(dst, src, n, info(src));
|
||||
case hipMemcpyDeviceToDevice: hsa_memory_copy(dst, src, n); break;
|
||||
default: return generic_copy(dst, src, n, info(dst), info(src));
|
||||
}
|
||||
}
|
||||
|
||||
hipError_t memcpyAsync(void* dst, const void* src, size_t sizeBytes,
|
||||
hipMemcpyKind kind, hipStream_t stream) {
|
||||
if (sizeBytes == 0) return hipSuccess;
|
||||
if (!dst || !src) return hipErrorInvalidValue;
|
||||
|
||||
if (!(stream = ihipSyncAndResolveStream(stream))) {
|
||||
return hipErrorInvalidValue;
|
||||
}
|
||||
|
||||
try {
|
||||
stream = ihipSyncAndResolveStream(stream);
|
||||
|
||||
if (!stream) return hipErrorInvalidValue;
|
||||
|
||||
stream->locked_copyAsync(dst, src, sizeBytes, kind);
|
||||
}
|
||||
catch (ihipException& ex) {
|
||||
e = ex._code;
|
||||
catch (const ihipException& ex) {
|
||||
return ex._code;
|
||||
}
|
||||
catch (const std::exception& ex) {
|
||||
std::cerr << ex.what() << std::endl;
|
||||
throw;
|
||||
}
|
||||
catch (...) {
|
||||
return hipErrorUnknown;
|
||||
}
|
||||
|
||||
return hipSuccess;
|
||||
}
|
||||
|
||||
hipError_t memcpySync(void* dst, const void* src, size_t sizeBytes,
|
||||
hipMemcpyKind kind, hipStream_t stream) {
|
||||
if (sizeBytes == 0) return hipSuccess;
|
||||
if (!dst || !src) return hipErrorInvalidValue;
|
||||
|
||||
try {
|
||||
stream = ihipSyncAndResolveStream(stream);
|
||||
|
||||
if (!stream) return hipErrorInvalidValue;
|
||||
|
||||
LockedAccessor_StreamCrit_t cs{stream->criticalData()};
|
||||
cs->_av.wait();
|
||||
|
||||
memcpy_impl(dst, src, sizeBytes, kind);
|
||||
cs->_last_op_was_a_copy = true;
|
||||
}
|
||||
catch (const ihipException& ex) {
|
||||
return ex._code;
|
||||
}
|
||||
catch (const std::exception& ex) {
|
||||
std::cerr << ex.what() << std::endl;
|
||||
throw;
|
||||
}
|
||||
catch (...) {
|
||||
return hipErrorUnknown;
|
||||
@@ -896,20 +1246,8 @@ hipError_t hipMemcpyToSymbol(void* dst, const void* src, size_t count,
|
||||
|
||||
tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbol_name, dst);
|
||||
|
||||
if (dst == nullptr) {
|
||||
return ihipLogStatus(hipErrorInvalidSymbol);
|
||||
}
|
||||
|
||||
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
|
||||
|
||||
if (kind == hipMemcpyHostToDevice || kind == hipMemcpyDefault ||
|
||||
kind == hipMemcpyDeviceToDevice || kind == hipMemcpyHostToHost) {
|
||||
stream->locked_copySync((char*)dst+offset, (void*)src, count, kind, false);
|
||||
} else {
|
||||
return ihipLogStatus(hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
return ihipLogStatus(hipSuccess);
|
||||
return ihipLogStatus(
|
||||
hipMemcpy(static_cast<char*>(dst) + offset, src, count, kind));
|
||||
}
|
||||
|
||||
hipError_t hipMemcpyFromSymbol(void* dst, const void* src, size_t count,
|
||||
@@ -920,20 +1258,8 @@ hipError_t hipMemcpyFromSymbol(void* dst, const void* src, size_t count,
|
||||
|
||||
tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbol_name, dst);
|
||||
|
||||
if (dst == nullptr) {
|
||||
return ihipLogStatus(hipErrorInvalidSymbol);
|
||||
}
|
||||
|
||||
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
|
||||
|
||||
if (kind == hipMemcpyDefault || kind == hipMemcpyDeviceToHost ||
|
||||
kind == hipMemcpyDeviceToDevice || kind == hipMemcpyHostToHost) {
|
||||
stream->locked_copySync((void*)dst, (char*)src+offset, count, kind, false);
|
||||
} else {
|
||||
return ihipLogStatus(hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
return ihipLogStatus(hipSuccess);
|
||||
return ihipLogStatus(
|
||||
hipMemcpy(dst, static_cast<const char*>(src) + offset, count, kind));
|
||||
}
|
||||
|
||||
|
||||
@@ -995,120 +1321,49 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* src, size_t count,
|
||||
hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind) {
|
||||
HIP_INIT_SPECIAL_API(hipMemcpy, (TRACE_MCMD), dst, src, sizeBytes, kind);
|
||||
|
||||
hipError_t e = hipSuccess;
|
||||
|
||||
// Return success if number of bytes to copy is 0
|
||||
if (sizeBytes == 0) return ihipLogStatus(e);
|
||||
|
||||
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
|
||||
|
||||
hc::completion_future marker;
|
||||
|
||||
if(dst==NULL || src==NULL)
|
||||
{
|
||||
e=hipErrorInvalidValue;
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
try {
|
||||
stream->locked_copySync(dst, src, sizeBytes, kind);
|
||||
} catch (ihipException& ex) {
|
||||
e = ex._code;
|
||||
}
|
||||
|
||||
return ihipLogStatus(e);
|
||||
return ihipLogStatus(hip_internal::memcpySync(dst, src, sizeBytes, kind,
|
||||
hipStreamNull));
|
||||
}
|
||||
|
||||
|
||||
hipError_t hipMemcpyHtoD(hipDeviceptr_t dst, void* src, size_t sizeBytes) {
|
||||
HIP_INIT_SPECIAL_API(hipMemcpyHtoD, (TRACE_MCMD), dst, src, sizeBytes);
|
||||
|
||||
hipError_t e = hipSuccess;
|
||||
if (sizeBytes == 0) return ihipLogStatus(e);
|
||||
|
||||
if(dst==NULL || src==NULL){
|
||||
return ihipLogStatus(hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
|
||||
|
||||
hc::completion_future marker;
|
||||
|
||||
try {
|
||||
stream->locked_copySync((void*)dst, (void*)src, sizeBytes, hipMemcpyHostToDevice, false);
|
||||
} catch (ihipException& ex) {
|
||||
e = ex._code;
|
||||
}
|
||||
|
||||
return ihipLogStatus(e);
|
||||
return ihipLogStatus(hip_internal::memcpySync(dst, src, sizeBytes,
|
||||
hipMemcpyHostToDevice,
|
||||
hipStreamNull));
|
||||
}
|
||||
|
||||
|
||||
hipError_t hipMemcpyDtoH(void* dst, hipDeviceptr_t src, size_t sizeBytes) {
|
||||
HIP_INIT_SPECIAL_API(hipMemcpyDtoH, (TRACE_MCMD), dst, src, sizeBytes);
|
||||
|
||||
hipError_t e = hipSuccess;
|
||||
if (sizeBytes == 0) return ihipLogStatus(e);
|
||||
|
||||
if(dst==NULL || src==NULL){
|
||||
return ihipLogStatus(hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
|
||||
|
||||
hc::completion_future marker;
|
||||
|
||||
try {
|
||||
stream->locked_copySync((void*)dst, (void*)src, sizeBytes, hipMemcpyDeviceToHost, false);
|
||||
} catch (ihipException& ex) {
|
||||
e = ex._code;
|
||||
}
|
||||
|
||||
return ihipLogStatus(e);
|
||||
return ihipLogStatus(hip_internal::memcpySync(dst, src, sizeBytes,
|
||||
hipMemcpyDeviceToHost,
|
||||
hipStreamNull));
|
||||
}
|
||||
|
||||
hipError_t hipMemcpyDtoD(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes) {
|
||||
HIP_INIT_SPECIAL_API(hipMemcpyDtoD, (TRACE_MCMD), dst, src, sizeBytes);
|
||||
|
||||
hipError_t e = hipSuccess;
|
||||
if (sizeBytes == 0) return ihipLogStatus(e);
|
||||
|
||||
if(dst==NULL || src==NULL){
|
||||
return ihipLogStatus(hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
|
||||
|
||||
hc::completion_future marker;
|
||||
|
||||
try {
|
||||
stream->locked_copySync((void*)dst, (void*)src, sizeBytes, hipMemcpyDeviceToDevice, false);
|
||||
} catch (ihipException& ex) {
|
||||
e = ex._code;
|
||||
}
|
||||
|
||||
return ihipLogStatus(e);
|
||||
return ihipLogStatus(hip_internal::memcpySync(dst, src, sizeBytes,
|
||||
hipMemcpyDeviceToDevice,
|
||||
hipStreamNull));
|
||||
}
|
||||
|
||||
hipError_t hipMemcpyHtoH(void* dst, void* src, size_t sizeBytes) {
|
||||
HIP_INIT_SPECIAL_API(hipMemcpyHtoH, (TRACE_MCMD), dst, src, sizeBytes);
|
||||
|
||||
hipError_t e = hipSuccess;
|
||||
if (sizeBytes == 0) return ihipLogStatus(e);
|
||||
return ihipLogStatus(hip_internal::memcpySync(dst, src, sizeBytes,
|
||||
hipMemcpyHostToHost,
|
||||
hipStreamNull));
|
||||
}
|
||||
|
||||
if(dst==NULL || src==NULL){
|
||||
return ihipLogStatus(hipErrorInvalidValue);
|
||||
}
|
||||
hipError_t hipMemcpyWithStream(void* dst, const void* src, size_t sizeBytes,
|
||||
hipMemcpyKind kind, hipStream_t stream) {
|
||||
HIP_INIT_SPECIAL_API(hipMemcpyWithStream, (TRACE_MCMD), dst, src, sizeBytes,
|
||||
kind, stream);
|
||||
|
||||
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
|
||||
|
||||
hc::completion_future marker;
|
||||
try {
|
||||
stream->locked_copySync((void*)dst, (void*)src, sizeBytes, hipMemcpyHostToHost, false);
|
||||
} catch (ihipException& ex) {
|
||||
e = ex._code;
|
||||
}
|
||||
|
||||
return ihipLogStatus(e);
|
||||
return ihipLogStatus(hip_internal::memcpySync(dst, src, sizeBytes, kind,
|
||||
stream));
|
||||
}
|
||||
|
||||
hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind,
|
||||
|
||||
@@ -239,14 +239,7 @@ hipError_t ihipModuleLaunchKernel(TlsData *tls, hipFunction_t f, uint32_t global
|
||||
aql.header |= (1 << HSA_PACKET_HEADER_BARRIER);
|
||||
}
|
||||
|
||||
if (HCC_OPT_FLUSH) {
|
||||
aql.header |= (HSA_FENCE_SCOPE_AGENT << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) |
|
||||
(HSA_FENCE_SCOPE_AGENT << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE);
|
||||
} else {
|
||||
aql.header |= (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) |
|
||||
(HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE);
|
||||
};
|
||||
|
||||
aql.header |= lp.launch_fence;
|
||||
|
||||
hc::completion_future cf;
|
||||
|
||||
|
||||
@@ -0,0 +1,57 @@
|
||||
/*
|
||||
Copyright (c) 2015-present Advanced Micro Devices, Inc. All rights reserved.
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/*
|
||||
* Conformance test for checking functionality of
|
||||
* hipError_t hipMemcpyPeer(void* dst, int dstDeviceId, const void* src, int srcDeviceId, size_t
|
||||
* sizeBytes);
|
||||
*/
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../../test_common.cpp
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include "test_common.h"
|
||||
|
||||
int main() {
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
int numDevices = 0;
|
||||
int *A_d, *B_d, *C_d;
|
||||
int *A_h, *B_h, *C_h;
|
||||
|
||||
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
|
||||
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
|
||||
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
|
||||
HIPCHECK(hipMemcpyWithStream(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream));
|
||||
HIPCHECK(hipMemcpyWithStream(B_d, B_h, Nbytes, hipMemcpyHostToDevice,stream));
|
||||
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream,
|
||||
static_cast<const int*>(A_d), static_cast<const int*>(B_d), C_d, N);
|
||||
HIPCHECK(hipStreamSynchronize(stream));
|
||||
HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
|
||||
HipTest::checkVectorADD(A_h, B_h, C_h, N);
|
||||
|
||||
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
passed();
|
||||
}
|
||||
Reference in New Issue
Block a user