diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h index 2cf14156a2..2c75b584c4 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -1159,7 +1159,7 @@ hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t siz */ hipError_t hipMemcpyToSymbolAsync(const char* symbolName, const void *src, size_t sizeBytes, size_t offset, hipMemcpyKind kind, hipStream_t stream); - +hipError_t hipMemcpyFromSymbolAsync(void *dst, const char* symbolName, size_t sizeBytes, size_t offset, hipMemcpyKind kind, hipStream_t stream); /** * @brief Copy data from src to dst asynchronously. diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index 68bf74cd21..a119ea1c54 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/projects/clr/hipamd/src/hip_hcc.cpp @@ -1780,7 +1780,7 @@ 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) +void ihipStream_t::lockedSymbolCopySync(hc::accelerator &acc, void* dst, void* src, size_t sizeBytes, unsigned kind) { if(kind == hipMemcpyHostToHost){ acc.memcpy_symbol(dst, (void*)src, sizeBytes, Kalmar::hcMemcpyHostToHost); @@ -1796,11 +1796,18 @@ void ihipStream_t::lockedSymbolCopySync(hc::accelerator &acc, void* dst, const v } } -void ihipStream_t::lockedSymbolCopyAsync(hc::accelerator &acc, void* dst, const void* src, size_t sizeBytes, unsigned kind) +void ihipStream_t::lockedSymbolCopyAsync(hc::accelerator &acc, void* dst, 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); + if(kind == hipMemcpyHostToDevice) { + hc::AmPointerInfo dstPtrInfo(NULL, dst, sizeBytes, acc, true, false); + hc::am_memtracker_add(dst, dstPtrInfo); + locked_getAv()->copy_async((void*)src, dst, sizeBytes); + } + if(kind == hipMemcpyDeviceToHost) { + hc::AmPointerInfo srcPtrInfo(NULL, src, sizeBytes, acc, true, false); + hc::am_memtracker_add(src, srcPtrInfo); + locked_getAv()->copy_async((void*)src, dst, sizeBytes); + } } void ihipStream_t::locked_copyAsync(void* dst, const void* src, size_t sizeBytes, unsigned kind) @@ -1903,7 +1910,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 5509e4aa10..d7d92a221c 100644 --- a/projects/clr/hipamd/src/hip_hcc.h +++ b/projects/clr/hipamd/src/hip_hcc.h @@ -494,8 +494,8 @@ public: 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); + void lockedSymbolCopySync(hc::accelerator &acc, void *dst, void* src, size_t sizeBytes, unsigned kind); + void lockedSymbolCopyAsync(hc::accelerator &acc, void *dst, 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 df776aa707..cb265159ba 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -461,7 +461,7 @@ hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t cou if(kind == hipMemcpyHostToDevice || kind == hipMemcpyDeviceToHost || kind == hipMemcpyDeviceToDevice || kind == hipMemcpyHostToHost) { - stream->lockedSymbolCopySync(acc, dst, src, count + offset, kind); + stream->lockedSymbolCopySync(acc, dst, (void*)src, count + offset, kind); // acc.memcpy_symbol(dst, (void*)src, count+offset); } else { return ihipLogStatus(hipErrorInvalidValue); @@ -493,6 +493,44 @@ hipError_t hipMemcpyToSymbolAsync(const char* symbolName, const void *src, size_ return ihipLogStatus(hipErrorInvalidSymbol); } + if (stream) { + try { + stream->lockedSymbolCopyAsync(acc, dst, (void*)src, count + offset, kind); + } + catch (ihipException ex) { + e = ex._code; + } + } else { + e = hipErrorInvalidValue; + } + + return ihipLogStatus(e); +} + + +hipError_t hipMemcpyFromSymbolAsync(void* dst, const char* symbolName, size_t count, size_t offset, hipMemcpyKind kind, hipStream_t stream) +{ + HIP_INIT_CMD_API(symbolName, dst, count, offset, kind, stream); + + if(symbolName == nullptr) + { + return ihipLogStatus(hipErrorInvalidSymbol); + } + + hipError_t e = hipSuccess; + + auto ctx = ihipGetTlsDefaultCtx(); + + hc::accelerator acc = ctx->getDevice()->_acc; + + void *src = acc.get_symbol_address(symbolName); + tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbolName, src); + + if(src == nullptr || dst == nullptr) + { + return ihipLogStatus(hipErrorInvalidSymbol); + } + if (stream) { try { stream->lockedSymbolCopyAsync(acc, dst, src, count + offset, kind); @@ -505,8 +543,6 @@ hipError_t hipMemcpyToSymbolAsync(const char* symbolName, const void *src, size_ } return ihipLogStatus(e); - - } //--- diff --git a/projects/clr/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp b/projects/clr/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp index 49a719a8d4..476c5e0997 100644 --- a/projects/clr/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp +++ b/projects/clr/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp @@ -18,7 +18,7 @@ THE SOFTWARE. */ /* HIT_START - * BUILD: %t %s + * BUILD: %t %s * RUN: %t * HIT_END */ @@ -32,44 +32,53 @@ THE SOFTWARE. #define SIZE 1024*4 #ifdef __HIP_PLATFORM_HCC__ -__attribute__((address_space(1))) int global[NUM]; +__attribute__((address_space(1))) int globalIn[NUM]; +__attribute__((address_space(1))) int globalOut[NUM]; #endif #ifdef __HIP_PLATFORM_NVCC__ -__device__ int global[NUM]; +__device__ int globalIn[NUM]; +__device__ int globalOut[NUM]; #endif __global__ void Assign(hipLaunchParm lp, int* Out) { int tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; - Out[tid] = global[tid]; + Out[tid] = globalIn[tid]; + globalOut[tid] = globalIn[tid]; } int main() { - int *A, *Am, *B, *Ad; + int *A, *Am, *B, *Ad, *C, *Cm; A = new int[NUM]; B = new int[NUM]; + C = new int[NUM]; for(unsigned i=0;i