Merge "Add hipDrvMemcpy3D." into amd-master-next

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;

include/hip/hcc_detail/driver_types.h

@@ -263,26 +263,29 @@ typedef struct hipMemcpy3DParms {
 } hipMemcpy3DParms;
 
 typedef struct HIP_MEMCPY3D {
-    size_t Depth;
+  unsigned int srcXInBytes;
-    size_t Height;
+  unsigned int srcY;
-    size_t WidthInBytes;
+  unsigned int srcZ;
-    hipDeviceptr_t dstDevice;
+  unsigned int srcLOD;
-    size_t dstHeight;
+  hipMemoryType srcMemoryType;
-    void* dstHost;
+  const void* srcHost;
-    size_t dstLOD;
+  hipDeviceptr_t srcDevice;
-    hipMemoryType dstMemoryType;
+  hipArray_t srcArray;
-    size_t dstPitch;
+  unsigned int srcPitch;
-    size_t dstXInBytes;
+  unsigned int srcHeight;
-    size_t dstY;
+  unsigned int dstXInBytes;
-    size_t dstZ;
+  unsigned int dstY;
-    void* reserved0;
+  unsigned int dstZ;
-    void* reserved1;
+  unsigned int dstLOD;
-    hipDeviceptr_t srcDevice;
+  hipMemoryType dstMemoryType;
-    size_t srcHeight;
+  void* dstHost;
-    const void* srcHost;
+  hipDeviceptr_t dstDevice;
-    size_t srcLOD;
+  hipArray_t dstArray;
-    hipMemoryType srcMemoryType;
+  unsigned int dstPitch;
-    size_t srcPitch;
+  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,

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
 /**
  * @}

vdi/hip_conversions.hpp

@@ -25,36 +25,6 @@ THE SOFTWARE.
 #include <hip/hcc_detail/driver_types.h>
 #include <hip/hcc_detail/texture_types.h>
 
-// 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<hipMemoryType, hipMemoryType> getMemoryType(const hipMemcpyKind kind)
 }
 
 inline
-_HIP_MEMCPY3D getDrvMemcpy3DDesc(const hip_Memcpy2D& desc2D) {
+HIP_MEMCPY3D getDrvMemcpy3DDesc(const hip_Memcpy2D& desc2D) {
-  _HIP_MEMCPY3D desc3D = {};
+  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 getDrvMemcpy3DDesc(const hipMemcpy3DParms& desc) {
-  _HIP_MEMCPY3D descDrv = {};
+  HIP_MEMCPY3D descDrv = {};
 
   descDrv.WidthInBytes = desc.extent.width;
   descDrv.Height = desc.extent.height;

vdi/hip_hcc.def.in

@@ -89,6 +89,8 @@ hipMemcpy2DAsync
 hipMemcpy2DToArray
 hipMemcpy3D
 hipMemcpy3DAsync
+hipDrvMemcpy3D
+hipDrvMemcpy3DAsync
 hipMemcpyAsync
 hipMemcpyDtoD
 hipMemcpyDtoDAsync

vdi/hip_hcc.map.in

@@ -90,6 +90,8 @@ global:
     hipMemcpy2DToArray;
     hipMemcpy3D;
     hipMemcpy3DAsync;
+    hipDrvMemcpy3D;
+    hipDrvMemcpy3DAsync;
     hipMemcpyAsync;
     hipMemcpyDtoD;
     hipMemcpyDtoDAsync;

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));