From 87178f94563be66ebaa9dddc6ff77f29aec1a77b Mon Sep 17 00:00:00 2001 From: Vladislav Sytchenko Date: Sun, 1 Mar 2020 13:40:14 -0500 Subject: [PATCH] Fix hipMemcpy3d (partially) Incoming changes from upstream split the struct hipMemcpy3DParms into two separate ones - hipMemcpy3DParms and HIP_MEMCPY3D, which are cudaMemcpy3DParms and CUDA_MEMCPY3D equivalents respectively. Note that HIP_MEMCPY3D is missing half the members of CUDA_MEMCPY3D (this should be fixed in PR#1887). Work around this by using a substitute _HIP_MEMCPY3D struct for now. Change-Id: Ic15134e6deb260189b662b3804d2309a9b8473e9 --- vdi/hip_conversions.hpp | 34 ++++++++++++++- vdi/hip_memory.cpp | 91 ++++++++++++++++++++++++++--------------- 2 files changed, 91 insertions(+), 34 deletions(-) diff --git a/vdi/hip_conversions.hpp b/vdi/hip_conversions.hpp index 321ef7e024..b34680ec72 100644 --- a/vdi/hip_conversions.hpp +++ b/vdi/hip_conversions.hpp @@ -25,6 +25,36 @@ THE SOFTWARE. #include #include +// HIP_MEMCPY3D is currently broken. +// TODO remove this struct once the headers will be fixed. +struct _HIP_MEMCPY3D { + unsigned int srcXInBytes; + unsigned int srcY; + unsigned int srcZ; + unsigned int srcLOD; + hipMemoryType srcMemoryType; + const void* srcHost; + hipDeviceptr_t srcDevice; + hipArray_t srcArray; + unsigned int srcPitch; + unsigned int srcHeight; + + unsigned int dstXInBytes; + unsigned int dstY; + unsigned int dstZ; + unsigned int dstLOD; + hipMemoryType dstMemoryType; + void* dstHost; + hipDeviceptr_t dstDevice; + hipArray_t dstArray; + unsigned int dstPitch; + unsigned int dstHeight; + + unsigned int WidthInBytes; + unsigned int Height; + unsigned int Depth; +}; + namespace hip { inline @@ -549,8 +579,8 @@ std::pair getMemoryType(const hipMemcpyKind kind) } inline -hipMemcpy3DParms getMemcpy3DParms(const hip_Memcpy2D& desc2D) { - hipMemcpy3DParms desc3D = {}; +_HIP_MEMCPY3D getMemcpy3DParms(const hip_Memcpy2D& desc2D) { + _HIP_MEMCPY3D desc3D = {}; desc3D.srcXInBytes = desc2D.srcXInBytes; desc3D.srcY = desc2D.srcY; diff --git a/vdi/hip_memory.cpp b/vdi/hip_memory.cpp index 6d68cfd9f5..7b4fece3b2 100644 --- a/vdi/hip_memory.cpp +++ b/vdi/hip_memory.cpp @@ -1324,7 +1324,7 @@ hipError_t ihipMemcpyAtoH(hipArray* srcArray, return hipSuccess; } -hipError_t ihipMemcpyParam3D(const hipMemcpy3DParms* pCopy, +hipError_t ihipMemcpyParam3D(const _HIP_MEMCPY3D* pCopy, hipStream_t stream, bool isAsync = false) { // If {src/dst}MemoryType is hipMemoryTypeUnified, {src/dst}Device and {src/dst}Pitch specify the (unified virtual address space) @@ -1379,13 +1379,13 @@ hipError_t ihipMemcpyParam3D(const hipMemcpy3DParms* pCopy, hipError_t ihipMemcpyParam2D(const hip_Memcpy2D* pCopy, hipStream_t stream, bool isAsync = false) { - hipMemcpy3DParms desc = hip::getMemcpy3DParms(*pCopy); + _HIP_MEMCPY3D desc = hip::getMemcpy3DParms(*pCopy); return ihipMemcpyParam3D(&desc, stream, isAsync); } hipError_t ihipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, - size_t height, hipMemcpyKind kind, hipStream_t stream, bool isAsync = false) { + size_t height, hipMemcpyKind kind, hipStream_t stream, bool isAsync = false) { hip_Memcpy2D desc = {}; desc.srcXInBytes = 0; @@ -1537,55 +1537,70 @@ hipError_t hipMemcpyAtoH(void* dstHost, hipError_t ihipMemcpy3D(const hipMemcpy3DParms* p, hipStream_t stream, bool isAsync = false) { - // Having src and dst be an array is ambigous, since we can't tell - // if the user intended to call runtime or driver version of hipMemcpy3D(). - // For now hope that we never encounter this case. - assert((p->srcArray == nullptr) || (p->dstArray == nullptr)); + // The struct passed to hipMemcpy3D() must specify one of srcArray or srcPtr and one of dstArray or dstPtr. + // Passing more than one non-zero source or destination will cause hipMemcpy3D() to return an error. + if (((p->srcArray != nullptr) && (p->srcPtr.ptr != nullptr)) || + ((p->dstArray != nullptr) && (p->dstPtr.ptr != nullptr))) { + return hipErrorInvalidValue; + } - // Now we need to patch the user provided struct if they intended on calling the runtime version. - hipMemcpy3DParms pCopy = {}; - std::memcpy(&pCopy, p, sizeof(hipMemcpy3DParms)); + // If the source and destination are both arrays, hipMemcpy3D() will return an error if they do not have the same element size. + if (((p->srcArray != nullptr) && (p->dstArray != nullptr)) && + (hip::getElementSize(p->dstArray->Format) != hip::getElementSize(p->dstArray->Format))) { + return hipErrorInvalidValue; + } - if ((p->srcPtr.ptr != nullptr) || (p->dstPtr.ptr != nullptr)) { - pCopy.WidthInBytes = p->extent.width; - pCopy.Height = p->extent.height; - pCopy.Depth = p->extent.depth; + _HIP_MEMCPY3D pCopy = {}; + + pCopy.WidthInBytes = p->extent.width; + pCopy.Height = p->extent.height; + pCopy.Depth = p->extent.depth; + + pCopy.srcXInBytes = p->srcPos.x; + pCopy.srcY = p->srcPos.y; + pCopy.srcZ = p->srcPos.z; + pCopy.srcLOD = 0; + + pCopy.dstXInBytes = p->dstPos.x; + pCopy.dstY = p->dstPos.y; + pCopy.dstZ = p->dstPos.z; + pCopy.dstLOD = 0; + + if (p->srcArray != nullptr) { + pCopy.srcMemoryType = hipMemoryTypeArray; + // When reffering to array memory, hipPos::x is in elements. + pCopy.srcXInBytes *= hip::getElementSize(p->dstArray->Format); } if (p->srcPtr.ptr != nullptr) { - pCopy.srcXInBytes = p->srcPos.x; - pCopy.srcY = p->srcPos.y; - pCopy.srcZ = p->srcPos.z; - pCopy.srcLOD = 0; pCopy.srcMemoryType = std::get<0>(hip::getMemoryType(p->kind)); pCopy.srcHost = p->srcPtr.ptr; pCopy.srcDevice = p->srcPtr.ptr; pCopy.srcPitch = p->srcPtr.pitch; pCopy.srcHeight = p->srcPtr.ysize; + } - if (p->dstArray != nullptr) { - pCopy.dstMemoryType = hipMemoryTypeArray; - // When reffering to array memory, hipExtent::width is in elements. - pCopy.WidthInBytes *= hip::getElementSize(p->dstArray->Format); - } + if (p->dstArray != nullptr) { + pCopy.dstMemoryType = hipMemoryTypeArray; + // When reffering to array memory, hipPos::x is in elements. + pCopy.srcXInBytes *= hip::getElementSize(p->dstArray->Format); } if (p->dstPtr.ptr != nullptr) { - pCopy.dstXInBytes = p->dstPos.x; - pCopy.dstY = p->dstPos.y; - pCopy.dstZ = p->dstPos.z; - pCopy.dstLOD = 0; pCopy.dstMemoryType = std::get<1>(hip::getMemoryType(p->kind)); pCopy.dstHost = p->dstPtr.ptr; pCopy.dstDevice = p->dstPtr.ptr; pCopy.dstPitch = p->dstPtr.pitch; pCopy.dstHeight = p->dstPtr.ysize; + } - if (p->srcArray != nullptr) { - pCopy.srcMemoryType = hipMemoryTypeArray; - // When reffering to array memory, hipExtent::width is in elements. - pCopy.WidthInBytes *= hip::getElementSize(p->srcArray->Format); - } + // If a HIP array is participating in the copy, the extent is defined in terms of that array's elements. + if ((p->srcArray != nullptr) && (p->dstArray == nullptr)) { + pCopy.WidthInBytes *= hip::getElementSize(p->srcArray->Format); + } else if ((p->srcArray == nullptr) && (p->dstArray != nullptr)) { + pCopy.WidthInBytes *= hip::getElementSize(p->dstArray->Format); + } else if ((p->srcArray != nullptr) && (p->dstArray != nullptr)) { + pCopy.WidthInBytes *= hip::getElementSize(p->dstArray->Format); } return ihipMemcpyParam3D(&pCopy, stream, isAsync); @@ -1603,6 +1618,18 @@ hipError_t hipMemcpy3DAsync(const hipMemcpy3DParms* p, hipStream_t stream) { HIP_RETURN(ihipMemcpy3D(p, stream, true)); } +hipError_t hipDrvMemcpy3D(const _HIP_MEMCPY3D* pCopy) { + HIP_INIT_API(hipDrvMemcpy3D, pCopy); + + HIP_RETURN(ihipMemcpyParam3D(pCopy, nullptr)); +} + +hipError_t hipDrvMemcpy3DAsync(const _HIP_MEMCPY3D* pCopy, hipStream_t stream) { + HIP_INIT_API(hipDrvMemcpy3DAsync, pCopy, stream); + + HIP_RETURN(ihipMemcpyParam3D(pCopy, stream, true)); +} + hipError_t ihipMemset(void* dst, int value, size_t valueSize, size_t sizeBytes, hipStream_t stream, bool isAsync = false) { if (sizeBytes == 0) {