AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge pull request #484 from gargrahul/fix_malloc_hiphostreg

Browse Source 
Fix memcpy2D for malloc+ hostRegister

[ROCm/hip commit: df450c6680]
This commit is contained in:
Maneesh Gupta
2018-06-01 16:53:25 +05:30
committed by GitHub
parent 121993cae8 46e623fb31
commit eded0da7b5
1 changed files with 41 additions and 25 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/src/hip_memory.cpp
+41 -25
View File
@@ -1586,23 +1586,22 @@ hipError_t ihipMemset(void* dst, int value, size_t sizeBytes, hipStream_t strea
return e; return e;
}; };
int isLockedPointer(const void *ptr) hipError_t getLockedPointer(const void *hostPtr, size_t dataLen, void **devicePtrPtr)
{ {
hsa_amd_pointer_info_t info; hc::accelerator acc;
int isLocked = 0;
info.size = sizeof(info); #if (__hcc_workweek__ >= 17332)
hsa_status_t hsa_status = hsa_amd_pointer_info(const_cast<void*>(ptr), &info, nullptr, nullptr, nullptr); hc::AmPointerInfo amPointerInfo(NULL, NULL, NULL, 0, acc, 0, 0);
if(hsa_status != HSA_STATUS_SUCCESS) { #else
return -1; hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0);
} #endif
am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, hostPtr);
if((info.type == HSA_EXT_POINTER_TYPE_HSA) || (info.type == HSA_EXT_POINTER_TYPE_LOCKED)) { if (status == AM_SUCCESS) {
isLocked = 1; *devicePtrPtr = (char*)amPointerInfo._devicePointer;
} return(hipSuccess);
};
return isLocked; return(hipErrorHostMemoryNotRegistered);
} };
// TODO - review and optimize // TODO - review and optimize
hipError_t ihipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, hipError_t ihipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width,
@@ -1611,12 +1610,20 @@ hipError_t ihipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
int isLocked = 0; int isLocked = 0;
if(kind == hipMemcpyHostToDevice) { void *pinnedPtr=NULL;
isLocked = isLockedPointer(src); void *actualSrc = (void*)src;
void *actualDest = dst;
if(kind == hipMemcpyHostToDevice ) {
if(getLockedPointer((void*)src, spitch, &pinnedPtr) == hipSuccess ){
isLocked = 1;
actualSrc = pinnedPtr;
}
} else if(kind == hipMemcpyDeviceToHost) { } else if(kind == hipMemcpyDeviceToHost) {
isLocked = isLockedPointer(dst); if(getLockedPointer((void*)dst, dpitch, &pinnedPtr) == hipSuccess ){
isLocked = 1;
actualDest = pinnedPtr;
}
} }
hc::completion_future marker; hc::completion_future marker;
hipError_t e = hipSuccess; hipError_t e = hipSuccess;
@@ -1624,12 +1631,12 @@ hipError_t ihipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch
stream->locked_copySync((void*)dst, (void*)src, width*height, kind, false); stream->locked_copySync((void*)dst, (void*)src, width*height, kind, false);
} else { } else {
try { try {
if(isLocked) { if(!isLocked) {
for (int i = 0; i < height; ++i) for (int i = 0; i < height; ++i)
stream->locked_copySync((unsigned char*)dst + i * dpitch, stream->locked_copySync((unsigned char*)dst + i * dpitch,
(unsigned char*)src + i * spitch, width, kind); (unsigned char*)src + i * spitch, width, kind);
} else { } else {
ihipMemcpy2dKernel<uint32_t> (stream, static_cast<uint32_t*> (dst), static_cast<const uint32_t*> (src), width, height, dpitch, spitch); ihipMemcpy2dKernel<uint32_t> (stream, static_cast<uint32_t*> (actualDest), static_cast<const uint32_t*> (actualSrc), width, height, dpitch, spitch);
stream->locked_wait(); stream->locked_wait();
} }
} catch (ihipException& ex) { } catch (ihipException& ex) {
@@ -1654,10 +1661,19 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp
if (dst == nullptr || src == nullptr || width > dpitch || width > spitch) return ihipLogStatus(hipErrorInvalidValue); if (dst == nullptr || src == nullptr || width > dpitch || width > spitch) return ihipLogStatus(hipErrorInvalidValue);
hipError_t e = hipSuccess; hipError_t e = hipSuccess;
int isLocked = 0; int isLocked = 0;
if(kind == hipMemcpyHostToDevice) { void *pinnedPtr=NULL;
isLocked = isLockedPointer(src); void *actualSrc = (void*)src;
void *actualDest = dst;
if(kind == hipMemcpyHostToDevice ) {
if(getLockedPointer((void*)src, spitch, &pinnedPtr) == hipSuccess ){
isLocked = 1;
actualSrc = pinnedPtr;
}
} else if(kind == hipMemcpyDeviceToHost) { } else if(kind == hipMemcpyDeviceToHost) {
isLocked = isLockedPointer(dst); if(getLockedPointer((void*)dst, dpitch, &pinnedPtr) == hipSuccess ){
isLocked = 1;
actualDest = pinnedPtr;
}
} }
if((width == dpitch) && (width == spitch)) { if((width == dpitch) && (width == spitch)) {
hip_internal::memcpyAsync(dst, src, width*height, kind, stream); hip_internal::memcpyAsync(dst, src, width*height, kind, stream);
@@ -1668,7 +1684,7 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp
e = hip_internal::memcpyAsync((unsigned char*)dst + i * dpitch, e = hip_internal::memcpyAsync((unsigned char*)dst + i * dpitch,
(unsigned char*)src + i * spitch, width, kind, stream); (unsigned char*)src + i * spitch, width, kind, stream);
} else{ } else{
ihipMemcpy2dKernel<uint32_t> (stream, static_cast<uint32_t*> (dst), static_cast<const uint32_t*> (src), width, height, dpitch, spitch); ihipMemcpy2dKernel<uint32_t> (stream, static_cast<uint32_t*> (actualDest), static_cast<const uint32_t*> (actualSrc), width, height, dpitch, spitch);
} }
} catch (ihipException& ex) { } catch (ihipException& ex) {
e = ex._code; e = ex._code;