From ecd7c99b49160866bba4d17501de52c3ac78728b Mon Sep 17 00:00:00 2001 From: Vladislav Sytchenko Date: Mon, 9 Mar 2020 14:05:16 -0400 Subject: [PATCH] Add hipDrvMemcpy3D. This is the equivalent of cuMemcpy3D. Change-Id: Ib2e06dbd6f5093c931cdfd36c87617f32acffc2d --- bin/hipify-perl | 6 ++++ include/hip/hcc_detail/driver_types.h | 43 +++++++++++++----------- include/hip/hcc_detail/hip_runtime_api.h | 25 ++++++++++++++ vdi/hip_conversions.hpp | 38 +++------------------ vdi/hip_hcc.def.in | 2 ++ vdi/hip_hcc.map.in | 2 ++ vdi/hip_memory.cpp | 10 +++--- 7 files changed, 67 insertions(+), 59 deletions(-) diff --git a/bin/hipify-perl b/bin/hipify-perl index f6de5abae4..62cd8436f8 100755 --- a/bin/hipify-perl +++ b/bin/hipify-perl @@ -223,6 +223,10 @@ sub simpleSubstitutions { $ft{'memory'} += s/\bcuMemcpy2DAsync\b/hipMemcpyParam2DAsync/g; $ft{'memory'} += s/\bcuMemcpy2DAsync_v2\b/hipMemcpyParam2DAsync/g; $ft{'memory'} += s/\bcuMemcpy2D_v2\b/hipMemcpyParam2D/g; + $ft{'memory'} += s/\bcuMemcpy3D\b/hipDrvMemcpy3D/g; + $ft{'memory'} += s/\bcuMemcpy3DAsync\b/hipDrvMemcpy3DAsync/g; + $ft{'memory'} += s/\bcuMemcpy3D_v2\b/hipDrvMemcpy3D/g; + $ft{'memory'} += s/\bcuMemcpy3DAsync_v2\b/hipDrvMemcpy3DAsync/g; $ft{'memory'} += s/\bcuMemcpyAtoH\b/hipMemcpyAtoH/g; $ft{'memory'} += s/\bcuMemcpyAtoH_v2\b/hipMemcpyAtoH/g; $ft{'memory'} += s/\bcuMemcpyDtoD\b/hipMemcpyDtoD/g; @@ -938,6 +942,8 @@ sub simpleSubstitutions { $ft{'type'} += s/\bCUDA_ARRAY_DESCRIPTOR_st\b/HIP_ARRAY_DESCRIPTOR/g; $ft{'type'} += s/\bCUDA_MEMCPY2D\b/hip_Memcpy2D/g; $ft{'type'} += s/\bCUDA_MEMCPY2D_st\b/hip_Memcpy2D/g; + $ft{'type'} += s/\bCUDA_MEMCPY3D\b/HIP_MEMCPY3D/g; + $ft{'type'} += s/\bCUDA_MEMCPY3D_st\b/HIP_MEMCPY3D/g; $ft{'type'} += s/\bCUaddress_mode\b/hipTextureAddressMode/g; $ft{'type'} += s/\bCUaddress_mode_enum\b/hipTextureAddressMode/g; $ft{'type'} += s/\bCUarray\b/hipArray */g; diff --git a/include/hip/hcc_detail/driver_types.h b/include/hip/hcc_detail/driver_types.h index 1941f44617..d24f097f84 100644 --- a/include/hip/hcc_detail/driver_types.h +++ b/include/hip/hcc_detail/driver_types.h @@ -263,26 +263,29 @@ typedef struct hipMemcpy3DParms { } hipMemcpy3DParms; typedef struct HIP_MEMCPY3D { - size_t Depth; - size_t Height; - size_t WidthInBytes; - hipDeviceptr_t dstDevice; - size_t dstHeight; - void* dstHost; - size_t dstLOD; - hipMemoryType dstMemoryType; - size_t dstPitch; - size_t dstXInBytes; - size_t dstY; - size_t dstZ; - void* reserved0; - void* reserved1; - hipDeviceptr_t srcDevice; - size_t srcHeight; - const void* srcHost; - size_t srcLOD; - hipMemoryType srcMemoryType; - size_t srcPitch; + 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; } HIP_MEMCPY3D; static inline struct hipPitchedPtr make_hipPitchedPtr(void* d, size_t p, size_t xsz, diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index 1f0b474863..2360333b1a 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -2165,6 +2165,31 @@ hipError_t hipMemcpy3D(const struct hipMemcpy3DParms* p); */ hipError_t hipMemcpy3DAsync(const struct hipMemcpy3DParms* p, hipStream_t stream __dparm(0)); +/** + * @brief Copies data between host and device. + * + * @param[in] pCopy 3D memory copy parameters + * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue, + * #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection + * + * @see hipMemcpy, hipMemcpy2DToArray, hipMemcpy2D, hipMemcpyFromArray, hipMemcpyToSymbol, + * hipMemcpyAsync + */ +hipError_t hipDrvMemcpy3D(const HIP_MEMCPY3D* pCopy); + +/** + * @brief Copies data between host and device asynchronously. + * + * @param[in] pCopy 3D memory copy parameters + * @param[in] stream Stream to use + * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue, + * #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection + * + * @see hipMemcpy, hipMemcpy2DToArray, hipMemcpy2D, hipMemcpyFromArray, hipMemcpyToSymbol, + * hipMemcpyAsync + */ +hipError_t hipDrvMemcpy3DAsync(const HIP_MEMCPY3D* pCopy, hipStream_t stream); + // doxygen end Memory /** * @} diff --git a/vdi/hip_conversions.hpp b/vdi/hip_conversions.hpp index ff88ae4828..c1ed166e16 100644 --- a/vdi/hip_conversions.hpp +++ b/vdi/hip_conversions.hpp @@ -25,36 +25,6 @@ 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 @@ -618,8 +588,8 @@ std::pair getMemoryType(const hipMemcpyKind kind) } inline -_HIP_MEMCPY3D getDrvMemcpy3DDesc(const hip_Memcpy2D& desc2D) { - _HIP_MEMCPY3D desc3D = {}; +HIP_MEMCPY3D getDrvMemcpy3DDesc(const hip_Memcpy2D& desc2D) { + HIP_MEMCPY3D desc3D = {}; desc3D.srcXInBytes = desc2D.srcXInBytes; desc3D.srcY = desc2D.srcY; @@ -651,8 +621,8 @@ _HIP_MEMCPY3D getDrvMemcpy3DDesc(const hip_Memcpy2D& desc2D) { } inline -_HIP_MEMCPY3D getDrvMemcpy3DDesc(const hipMemcpy3DParms& desc) { - _HIP_MEMCPY3D descDrv = {}; +HIP_MEMCPY3D getDrvMemcpy3DDesc(const hipMemcpy3DParms& desc) { + HIP_MEMCPY3D descDrv = {}; descDrv.WidthInBytes = desc.extent.width; descDrv.Height = desc.extent.height; diff --git a/vdi/hip_hcc.def.in b/vdi/hip_hcc.def.in index 54492eb48e..4ec2f698c0 100644 --- a/vdi/hip_hcc.def.in +++ b/vdi/hip_hcc.def.in @@ -89,6 +89,8 @@ hipMemcpy2DAsync hipMemcpy2DToArray hipMemcpy3D hipMemcpy3DAsync +hipDrvMemcpy3D +hipDrvMemcpy3DAsync hipMemcpyAsync hipMemcpyDtoD hipMemcpyDtoDAsync diff --git a/vdi/hip_hcc.map.in b/vdi/hip_hcc.map.in index 03d48fe9e7..b8ee45d647 100644 --- a/vdi/hip_hcc.map.in +++ b/vdi/hip_hcc.map.in @@ -90,6 +90,8 @@ global: hipMemcpy2DToArray; hipMemcpy3D; hipMemcpy3DAsync; + hipDrvMemcpy3D; + hipDrvMemcpy3DAsync; hipMemcpyAsync; hipMemcpyDtoD; hipMemcpyDtoDAsync; diff --git a/vdi/hip_memory.cpp b/vdi/hip_memory.cpp index 84f8fa9def..55bc11512a 100644 --- a/vdi/hip_memory.cpp +++ b/vdi/hip_memory.cpp @@ -1332,7 +1332,7 @@ hipError_t ihipMemcpyAtoH(hipArray* srcArray, return hipSuccess; } -hipError_t ihipMemcpyParam3D(const _HIP_MEMCPY3D* 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) @@ -1387,7 +1387,7 @@ hipError_t ihipMemcpyParam3D(const _HIP_MEMCPY3D* pCopy, hipError_t ihipMemcpyParam2D(const hip_Memcpy2D* pCopy, hipStream_t stream, bool isAsync = false) { - _HIP_MEMCPY3D desc = hip::getDrvMemcpy3DDesc(*pCopy); + HIP_MEMCPY3D desc = hip::getDrvMemcpy3DDesc(*pCopy); return ihipMemcpyParam3D(&desc, stream, isAsync); } @@ -1558,7 +1558,7 @@ hipError_t ihipMemcpy3D(const hipMemcpy3DParms* p, return hipErrorInvalidValue; } - const _HIP_MEMCPY3D desc = hip::getDrvMemcpy3DDesc(*p); + const HIP_MEMCPY3D desc = hip::getDrvMemcpy3DDesc(*p); return ihipMemcpyParam3D(&desc, stream, isAsync); } @@ -1575,13 +1575,13 @@ hipError_t hipMemcpy3DAsync(const hipMemcpy3DParms* p, hipStream_t stream) { HIP_RETURN(ihipMemcpy3D(p, stream, true)); } -hipError_t hipDrvMemcpy3D(const _HIP_MEMCPY3D* pCopy) { +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) { +hipError_t hipDrvMemcpy3DAsync(const HIP_MEMCPY3D* pCopy, hipStream_t stream) { HIP_INIT_API(hipDrvMemcpy3DAsync, pCopy, stream); HIP_RETURN(ihipMemcpyParam3D(pCopy, stream, true));