From 0bf811b8758b5ca0068cf7569d50e111ccf77461 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Tue, 11 Oct 2016 13:29:46 -0500 Subject: [PATCH] added more changes to memcpytosymbol 1. Refactored code to use HCC internal APIs rather than HCC copy APIs 2. Added hipMemcpyToSymbolAsync 3. Added test for hipMemcpyToSymbolAsync 4. Added new error hipErrorInvalidSymbol Change-Id: I0e359b2d0ff5d682bbccdf9c2923e16b35e39497 --- .../include/hip/hcc_detail/hip_runtime_api.h | 21 +++++++ hipamd/include/hip/hip_runtime_api.h | 2 +- hipamd/src/hip_memory.cpp | 61 ++++++++++++++++--- .../src/deviceLib/hipTestDeviceSymbol.cpp | 20 +++++- 4 files changed, 92 insertions(+), 12 deletions(-) diff --git a/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/hipamd/include/hip/hcc_detail/hip_runtime_api.h index 1243934d01..4cbaf9ea4e 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -1062,6 +1062,27 @@ hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t siz hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t sizeBytes, size_t offset, hipMemcpyKind kind); +/** + * @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 + * + * The memory areas may not overlap. Symbol can either be a variable that resides in global or constant memory space, or it can be a character string, + * naming a variable that resides in global or constant memory space. Kind can be either hipMemcpyHostToDevice or hipMemcpyDeviceToDevice + * hipMemcpyToSymbolAsync() is asynchronous with respect to the host, so the call may return before copy is complete. + * TODO: cudaErrorInvalidSymbol and cudaErrorInvalidMemcpyDirection is not supported, use hipErrorUnknown for now. + * + * @param[in] symbolName - Symbol destination on device + * @param[in] src - Data being copy from + * @param[in] sizeBytes - Data size in bytes + * @param[in] offset - Offset from start of symbol in bytes + * @param[in] kind - Type of transfer + * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryFree, #hipErrorUnknown + * + * @see hipMemcpy, hipMemcpy2D, hipMemcpyToArray, hipMemcpy2DToArray, hipMemcpyFromArray, hipMemcpy2DFromArray, hipMemcpyArrayToArray, hipMemcpy2DArrayToArray, hipMemcpyFromSymbol, hipMemcpyAsync, hipMemcpy2DAsync, hipMemcpyToArrayAsync, hipMemcpy2DToArrayAsync, hipMemcpyFromArrayAsync, hipMemcpy2DFromArrayAsync, hipMemcpyToSymbolAsync, hipMemcpyFromSymbolAsync + */ +hipError_t hipMemcpyToSymbolAsync(const char* symbolName, const void *src, size_t sizeBytes, size_t offset, hipMemcpyKind kind, hipStream_t stream); + + + /** * @brief Copy data from src to dst asynchronously. * diff --git a/hipamd/include/hip/hip_runtime_api.h b/hipamd/include/hip/hip_runtime_api.h index 3406bcbbc9..f08f03e6f5 100644 --- a/hipamd/include/hip/hip_runtime_api.h +++ b/hipamd/include/hip/hip_runtime_api.h @@ -183,7 +183,7 @@ typedef enum hipError_t { hipErrorInvalidHandle = 400, hipErrorNotFound = 500, hipErrorIllegalAddress = 700, - + hipErrorInvalidSymbol = 701, // Runtime Error Codes start here. hipErrorMissingConfiguration = 1001, hipErrorMemoryAllocation = 1002, ///< Memory allocation error. diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index 278502973a..a7c19e8949 100644 --- a/hipamd/src/hip_memory.cpp +++ b/hipamd/src/hip_memory.cpp @@ -385,24 +385,67 @@ hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t cou { HIP_INIT_API(symbolName, src, count, offset, kind); -#ifdef USE_MEMCPYTOSYMBOL - if(kind != hipMemcpyHostToDevice) + if(symbolName == nullptr) { - return ihipLogStatus(hipErrorInvalidValue); + return ihipLogStatus(hipErrorInvalidSymbol); } - auto ctx = ihipGetTlsDefaultCtx(); - //hsa_signal_t depSignal; - //int depSignalCnt = ctx._default_stream->preCopyCommand(NULL, &depSignal, ihipCommandCopyH2D); - assert(0); // Need to properly synchronize the copy - do something with depSignal if != NULL. + auto ctx = ihipGetTlsDefaultCtx(); hc::accelerator acc = ctx->getDevice()->_acc; - acc.memcpy_symbol(symbolName, (void*) src,count, offset); -#endif + void *ptr = acc.get_symbol_address(symbolName); + + if(ptr == nullptr) + { + return ihipLogStatus(hipErrorInvalidSymbol); + } + + hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); + + stream->locked_copySync(ptr, src, count + offset, kind); + return ihipLogStatus(hipSuccess); } +hipError_t hipMemcpyToSymbolAsync(const char* symbolName, const void *src, size_t count, size_t offset, hipMemcpyKind kind, hipStream_t stream) +{ + HIP_INIT_API(symbolName, src, count, offset, kind, stream); + + if(symbolName == nullptr) + { + return ihipLogStatus(hipErrorInvalidSymbol); + } + + hipError_t e = hipSuccess; + + auto ctx = ihipGetTlsDefaultCtx(); + + hc::accelerator acc = ctx->getDevice()->_acc; + + void *ptr = acc.get_symbol_address(symbolName); + + if(ptr == nullptr) + { + return ihipLogStatus(hipErrorInvalidSymbol); + } + + if (stream) { + try { + stream->locked_copyAsync(ptr, src, count + offset, kind); + } + catch (ihipException ex) { + e = ex._code; + } + } else { + e = hipErrorInvalidValue; + } + + return ihipLogStatus(e); + + +} + //--- hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind) { diff --git a/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp b/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp index 359bc5d6db..6a71295bd2 100644 --- a/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp +++ b/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp @@ -38,13 +38,29 @@ int main() int *A, *B, *Ad; A = new int[NUM]; B = new int[NUM]; - for(unsigned i=0;i