From 150e690a3ab552faa458f4d2c944534873cea15c Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Mon, 23 Dec 2019 15:41:24 +0200 Subject: [PATCH] Fix late-coming issues. (#1724) Implementation for hipMemcpyWithStream. --- hipamd/include/hip/hcc_detail/grid_launch.h | 2 - .../include/hip/hcc_detail/hip_runtime_api.h | 3 + .../include/hip/nvcc_detail/hip_runtime_api.h | 12 + hipamd/src/hip_hcc.cpp | 18 +- hipamd/src/hip_hcc_internal.h | 10 +- hipamd/src/hip_memory.cpp | 517 +++++++++++++----- hipamd/src/hip_module.cpp | 9 +- .../runtimeApi/memory/hipMemcpyWithStream.cpp | 57 ++ 8 files changed, 473 insertions(+), 155 deletions(-) create mode 100644 hipamd/tests/src/runtimeApi/memory/hipMemcpyWithStream.cpp diff --git a/hipamd/include/hip/hcc_detail/grid_launch.h b/hipamd/include/hip/hcc_detail/grid_launch.h index 61fd9bdbe7..22841a5657 100644 --- a/hipamd/include/hip/hcc_detail/grid_launch.h +++ b/hipamd/include/hip/hcc_detail/grid_launch.h @@ -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. diff --git a/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/hipamd/include/hip/hcc_detail/hip_runtime_api.h index c263f605da..c619cfcc44 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -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 * diff --git a/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/hipamd/include/hip/nvcc_detail/hip_runtime_api.h index 9cb59f14ea..e8d7885a8c 100644 --- a/hipamd/include/hip/nvcc_detail/hip_runtime_api.h +++ b/hipamd/include/hip/nvcc_detail/hip_runtime_api.h @@ -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( diff --git a/hipamd/src/hip_hcc.cpp b/hipamd/src/hip_hcc.cpp index e93490a08e..360a68103e 100644 --- a/hipamd/src/hip_hcc.cpp +++ b/hipamd/src/hip_hcc.cpp @@ -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); diff --git a/hipamd/src/hip_hcc_internal.h b/hipamd/src/hip_hcc_internal.h index 08b30196b8..99dd3146b1 100644 --- a/hipamd/src/hip_hcc_internal.h +++ b/hipamd/src/hip_hcc_internal.h @@ -495,9 +495,10 @@ struct LockedBase { template class ihipStreamCriticalBase_t : public LockedBase { - 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 { return gotLock ? this : nullptr; }; - public: ihipStream_t* _parent; - hc::accelerator_view _av; - - private: + bool _last_op_was_a_copy; }; diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index 266f9b51d6..30258d7474 100644 --- a/hipamd/src/hip_memory.cpp +++ b/hipamd/src/hip_memory.cpp @@ -27,6 +27,7 @@ THE SOFTWARE. #include "hip_hcc_internal.h" #include "trace_helper.h" +#include #include __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 + 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(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 hsa{ + (hsa_init() == HSA_STATUS_SUCCESS) + ? reinterpret_cast(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(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(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 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(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(di.agentBaseAddress) + + (static_cast(dst) - + static_cast(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 lck{ + dst, [](void* p) { hsa_amd_memory_unlock(p); }}; + + throwing_result_check(hsa_amd_memory_lock(dst, n, &si.agentOwner, 1, + const_cast(&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(src))}; + + if (si.type == HSA_EXT_POINTER_TYPE_LOCKED) { + src = static_cast(si.agentBaseAddress) + + (static_cast(src) - + static_cast(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 lck{ + const_cast(src), [](void* p) { hsa_amd_memory_unlock(p); }}; + + throwing_result_check(hsa_amd_memory_lock(const_cast(src), n, + &di.agentOwner, 1, + const_cast(&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 lck0{ + nullptr, [](void* p) { hsa_amd_memory_unlock(p); }}; + std::unique_ptr 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(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(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, diff --git a/hipamd/src/hip_module.cpp b/hipamd/src/hip_module.cpp index 552ccfccac..d9fc79ce63 100644 --- a/hipamd/src/hip_module.cpp +++ b/hipamd/src/hip_module.cpp @@ -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; diff --git a/hipamd/tests/src/runtimeApi/memory/hipMemcpyWithStream.cpp b/hipamd/tests/src/runtimeApi/memory/hipMemcpyWithStream.cpp new file mode 100644 index 0000000000..d3a7158619 --- /dev/null +++ b/hipamd/tests/src/runtimeApi/memory/hipMemcpyWithStream.cpp @@ -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(A_d), static_cast(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(); +}