From fa734c5a6af7109b9862edf9757cbffcd6cf273b Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Fri, 26 Aug 2016 13:11:01 -0500 Subject: [PATCH] Added explicit memory copy direction apis - Fixed stale printf in context api - Added 4 sync memcpy apis 1. hipMemcpyHtoD 2. hipMemcpyDtoH 3. hipMemcpyDtoD 4. hipMemcpyHtoH - Added test for added apis Change-Id: I4a9c382445b62631f8d0bcbb9a670322288b72b1 [ROCm/clr commit: f722a132a4f61b251667b54a50537e52dfb1e1bb] --- .../clr/hipamd/include/hcc_detail/hip_hcc.h | 5 +- .../include/hcc_detail/hip_runtime_api.h | 8 ++ projects/clr/hipamd/src/hip_context.cpp | 1 - projects/clr/hipamd/src/hip_hcc.cpp | 10 +-- projects/clr/hipamd/src/hip_memory.cpp | 84 +++++++++++++++++++ .../clr/hipamd/tests/src/hipDrvMemcpy.cpp | 26 ++++++ 6 files changed, 125 insertions(+), 9 deletions(-) create mode 100644 projects/clr/hipamd/tests/src/hipDrvMemcpy.cpp diff --git a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h index 48f7dac318..449d9bc522 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h +++ b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h @@ -421,8 +421,9 @@ typedef uint64_t SeqNum_t ; ~ihipStream_t(); // kind is hipMemcpyKind - void copySync (LockedAccessor_StreamCrit_t &crit, void* dst, const void* src, size_t sizeBytes, unsigned kind); - void locked_copySync (void* dst, const void* src, size_t sizeBytes, unsigned kind); + void copySync (LockedAccessor_StreamCrit_t &crit, void* dst, const void* src, size_t sizeBytes, unsigned kind, bool resolveOn = true); + void locked_copySync (void* dst, const void* src, size_t sizeBytes, unsigned kind, bool resolveOn = true); + void copyAsync(void* dst, const void* src, size_t sizeBytes, unsigned kind); diff --git a/projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h index a2b5587ee6..7f79b64ac9 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h @@ -844,6 +844,14 @@ hipError_t hipHostFree(void* ptr); */ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind); +hipError_t hipMemcpyHtoD(hipDeviceptr dst, hipDeviceptr src, size_t sizeBytes); + +hipError_t hipMemcpyDtoH(hipDeviceptr dst, hipDeviceptr src, size_t sizeBytes); + +hipError_t hipMemcpyDtoD(hipDeviceptr dst, hipDeviceptr src, size_t sizeBytes); + +hipError_t hipMemcpyHtoH(hipDeviceptr dst, hipDeviceptr src, size_t sizeBytes); + /** * @brief Copies @p sizeBytes bytes from the memory area pointed to by @p src to the memory area pointed to by @p offset bytes from the start of symbol @p symbol. diff --git a/projects/clr/hipamd/src/hip_context.cpp b/projects/clr/hipamd/src/hip_context.cpp index c0c98de827..1caa3b4e42 100644 --- a/projects/clr/hipamd/src/hip_context.cpp +++ b/projects/clr/hipamd/src/hip_context.cpp @@ -32,7 +32,6 @@ thread_local std::stack tls_ctxStack; hipError_t ihipCtxStackUpdate() { //HIP_INIT_API(); - printf("Inside ihipCtxStackUpdate\n"); hipError_t e = hipSuccess; if(tls_ctxStack.empty()) { diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index 589b9ef47e..09cf46f8a7 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/projects/clr/hipamd/src/hip_hcc.cpp @@ -1563,7 +1563,7 @@ void ihipStream_t::setAsyncCopyAgents(unsigned kind, ihipCommand_t *commandType, } -void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const void* src, size_t sizeBytes, unsigned kind) +void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const void* src, size_t sizeBytes, unsigned kind, bool resolveOn) { ihipCtx_t *ctx = this->getCtx(); const ihipDevice_t *device = ctx->getDevice(); @@ -1582,7 +1582,7 @@ void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const bool dstInDeviceMem = dstPtrInfo._isInDeviceMem; // Resolve default to a specific Kind so we know which algorithm to use: - if (kind == hipMemcpyDefault) { + if (kind == hipMemcpyDefault && resolveOn) { kind = resolveMemcpyDirection(srcTracked, dstTracked, srcInDeviceMem, dstInDeviceMem); }; @@ -1753,14 +1753,12 @@ void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const // Sync copy that acquires lock: -void ihipStream_t::locked_copySync(void* dst, const void* src, size_t sizeBytes, unsigned kind) +void ihipStream_t::locked_copySync(void* dst, const void* src, size_t sizeBytes, unsigned kind, bool resolveOn) { LockedAccessor_StreamCrit_t crit (_criticalData); - copySync(crit, dst, src, sizeBytes, kind); + copySync(crit, dst, src, sizeBytes, kind, resolveOn); } - - void ihipStream_t::copyAsync(void* dst, const void* src, size_t sizeBytes, unsigned kind) { LockedAccessor_StreamCrit_t crit(_criticalData); diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index c0282dc372..50520a4e0e 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -446,12 +446,96 @@ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind e = ex._code; } + return ihipLogStatus(e); +} +hipError_t hipMemcpyHtoD(hipDeviceptr dst, hipDeviceptr src, size_t sizeBytes) +{ + HIP_INIT_API(dst, src, sizeBytes); + + hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); + + hc::completion_future marker; + + hipError_t e = hipSuccess; + + try { + + stream->locked_copySync((void*)dst, (void*)src, sizeBytes, hipMemcpyHostToDevice, false); + } + catch (ihipException ex) { + e = ex._code; + } return ihipLogStatus(e); } +hipError_t hipMemcpyDtoH(hipDeviceptr dst, hipDeviceptr src, size_t sizeBytes) +{ + HIP_INIT_API(dst, src, sizeBytes); + + hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); + + hc::completion_future marker; + + hipError_t e = hipSuccess; + + try { + + stream->locked_copySync((void*)dst, (void*)src, sizeBytes, hipMemcpyDeviceToHost, false); + } + catch (ihipException ex) { + e = ex._code; + } + + return ihipLogStatus(e); +} + +hipError_t hipMemcpyDtoD(hipDeviceptr dst, hipDeviceptr src, size_t sizeBytes) +{ + HIP_INIT_API(dst, src, sizeBytes); + + hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); + + hc::completion_future marker; + + hipError_t e = hipSuccess; + + try { + + stream->locked_copySync((void*)dst, (void*)src, sizeBytes, hipMemcpyDeviceToDevice, false); + } + catch (ihipException ex) { + e = ex._code; + } + + return ihipLogStatus(e); +} + +hipError_t hipMemcpyHtoH(hipDeviceptr dst, hipDeviceptr src, size_t sizeBytes) +{ + HIP_INIT_API(dst, src, sizeBytes); + + hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); + + hc::completion_future marker; + + hipError_t e = hipSuccess; + + try { + + stream->locked_copySync((void*)dst, (void*)src, sizeBytes, hipMemcpyHostToHost, false); + } + catch (ihipException ex) { + e = ex._code; + } + + return ihipLogStatus(e); +} + + + /** * @result #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidMemcpyDirection, * @result #hipErrorInvalidValue : If dst==NULL or src==NULL, or other bad argument. diff --git a/projects/clr/hipamd/tests/src/hipDrvMemcpy.cpp b/projects/clr/hipamd/tests/src/hipDrvMemcpy.cpp new file mode 100644 index 0000000000..55c9b18818 --- /dev/null +++ b/projects/clr/hipamd/tests/src/hipDrvMemcpy.cpp @@ -0,0 +1,26 @@ +#include +#include +#include + +#define LEN 1024 +#define SIZE LEN<<2 + +int main(){ + int *A, *B, *C; + hipDeviceptr Ad, Bd; + A = new int[LEN]; + B = new int[LEN]; + C = new int[LEN]; + for(int i=0;i