From f36992067b10af12cdab256d22cc3f53a19c7ff2 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Wed, 1 Feb 2017 17:54:59 -0600 Subject: [PATCH] fixed symbol memcpy issue Change-Id: I89d7401be51d194bcbf771020ba66e3d3b6a18f8 [ROCm/clr commit: 0d4e6ae60acd0646c66e3c5802239495490e2a5d] --- projects/clr/hipamd/src/hip_hcc.cpp | 23 +++++++++++++++++++ projects/clr/hipamd/src/hip_hcc.h | 8 +++---- projects/clr/hipamd/src/hip_memory.cpp | 22 +++++++++++------- .../src/deviceLib/hipTestDeviceSymbol.cpp | 8 +++++-- 4 files changed, 46 insertions(+), 15 deletions(-) diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index 20ff1722e3..68bf74cd21 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/projects/clr/hipamd/src/hip_hcc.cpp @@ -1780,6 +1780,28 @@ void ihipStream_t::locked_copySync(void* dst, const void* src, size_t sizeBytes, } } +void ihipStream_t::lockedSymbolCopySync(hc::accelerator &acc, void* dst, const void* src, size_t sizeBytes, unsigned kind) +{ + if(kind == hipMemcpyHostToHost){ + acc.memcpy_symbol(dst, (void*)src, sizeBytes, Kalmar::hcMemcpyHostToHost); + } + if(kind == hipMemcpyHostToDevice){ + acc.memcpy_symbol(dst, (void*)src, sizeBytes); + } + if(kind == hipMemcpyDeviceToDevice){ + acc.memcpy_symbol(dst, (void*)src, sizeBytes, Kalmar::hcMemcpyDeviceToDevice); + } + if(kind == hipMemcpyDeviceToHost){ + acc.memcpy_symbol(dst, (void*)src, sizeBytes, Kalmar::hcMemcpyDeviceToHost); + } +} + +void ihipStream_t::lockedSymbolCopyAsync(hc::accelerator &acc, void* dst, const void* src, size_t sizeBytes, unsigned kind) +{ + hc::AmPointerInfo dstPtrInfo(NULL, dst, sizeBytes, acc, true, false); + hc::am_memtracker_add(dst, dstPtrInfo); + locked_getAv()->copy_async((void*)src, dst, sizeBytes); +} void ihipStream_t::locked_copyAsync(void* dst, const void* src, size_t sizeBytes, unsigned kind) { @@ -1881,6 +1903,7 @@ void ihipStream_t::locked_copyAsync(void* dst, const void* src, size_t sizeBytes LockedAccessor_StreamCrit_t crit(_criticalData); this->ensureHaveQueue(crit); + #if USE_COPY_EXT_V2 crit->_av.copy_ext(src, dst, sizeBytes, hcCopyDir, srcPtrInfo, dstPtrInfo, copyDevice ? ©Device->getDevice()->_acc : nullptr, forceUnpinnedCopy); #else diff --git a/projects/clr/hipamd/src/hip_hcc.h b/projects/clr/hipamd/src/hip_hcc.h index 9998c4ec0e..28ae6a2e8e 100644 --- a/projects/clr/hipamd/src/hip_hcc.h +++ b/projects/clr/hipamd/src/hip_hcc.h @@ -273,7 +273,7 @@ extern void recordApiTrace(std::string *fullStr, const std::string &apiStr); API_TRACE(0, __VA_ARGS__); -// Like above, but will trace with DB_CMD. +// Like above, but will trace with DB_CMD. // Replace HIP_INIT_API with this call inside important APIs that launch work on the GPU: // kernel launches, copy commands, memory sets, etc. #define HIP_INIT_CMD_API(...) \ @@ -479,8 +479,6 @@ private: typedef ihipStreamCriticalBase_t ihipStreamCritical_t; typedef LockedAccessor LockedAccessor_StreamCrit_t; - - //--- // Internal stream structure. class ihipStream_t { @@ -494,10 +492,10 @@ public: // kind is hipMemcpyKind void locked_copySync (void* dst, const void* src, size_t sizeBytes, unsigned kind, bool resolveOn = true); - - void locked_copyAsync(void* dst, const void* src, size_t sizeBytes, unsigned kind); + void lockedSymbolCopySync(hc::accelerator &acc, void *dst, const void* src, size_t sizeBytes, unsigned kind); + void lockedSymbolCopyAsync(hc::accelerator &acc, void *dst, const void* src, size_t sizeBytes, unsigned kind); //--- // Member functions that begin with locked_ are thread-safe accessors - these acquire / release the critical mutex. diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index d66c151266..df776aa707 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -449,17 +449,23 @@ hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t cou hc::accelerator acc = ctx->getDevice()->_acc; - void *ptr = acc.get_symbol_address(symbolName); - tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbolName, ptr); + void *dst = acc.get_symbol_address(symbolName); + tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbolName, dst); - if(ptr == nullptr) + if(dst == nullptr) { return ihipLogStatus(hipErrorInvalidSymbol); } hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); - stream->locked_copySync(ptr, src, count + offset, kind); + if(kind == hipMemcpyHostToDevice || kind == hipMemcpyDeviceToHost || kind == hipMemcpyDeviceToDevice || kind == hipMemcpyHostToHost) + { + stream->lockedSymbolCopySync(acc, dst, src, count + offset, kind); + // acc.memcpy_symbol(dst, (void*)src, count+offset); + } else { + return ihipLogStatus(hipErrorInvalidValue); + } return ihipLogStatus(hipSuccess); } @@ -479,17 +485,17 @@ hipError_t hipMemcpyToSymbolAsync(const char* symbolName, const void *src, size_ hc::accelerator acc = ctx->getDevice()->_acc; - void *ptr = acc.get_symbol_address(symbolName); - tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbolName, ptr); + void *dst = acc.get_symbol_address(symbolName); + tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbolName, dst); - if(ptr == nullptr) + if(dst == nullptr) { return ihipLogStatus(hipErrorInvalidSymbol); } if (stream) { try { - stream->locked_copyAsync(ptr, src, count + offset, kind); + stream->lockedSymbolCopyAsync(acc, dst, src, count + offset, kind); } catch (ihipException ex) { e = ex._code; diff --git a/projects/clr/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp b/projects/clr/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp index f90ed38967..70fe794f67 100644 --- a/projects/clr/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp +++ b/projects/clr/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp @@ -47,7 +47,7 @@ __global__ void Assign(hipLaunchParm lp, int* Out) int main() { - int *A, *B, *Ad; + int *A, *Am, *B, *Ad; A = new int[NUM]; B = new int[NUM]; for(unsigned i=0;i