[HIP][HIPIFY] Make hipMemcpyParam2D coherent with cuMemcpy2D

+ Makes hip_Memcpy2D struct compatible with CUDA_MEMCPY2D struct
+ Add hipMemcpyParam2D support in nvcc fallback path
+ Update hipify-clang, tests and docs accordingly


[ROCm/hip commit: 8f059b0ee9]

projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md

@@ -861,7 +861,7 @@
 | `cuMemAllocManaged`                                       |                               |
 | `cuMemAllocPitch`                                         |                               |
 | `cuMemcpy`                                                |                               |
-| `cuMemcpy2D`                                              |                               |
+| `cuMemcpy2D`                                              | `hipMemcpyParam2D`            |
 | `cuMemcpy2DAsync`                                         |                               |
 | `cuMemcpy2DUnaligned`                                     |                               |
 | `cuMemcpy3D`                                              |                               |

projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp

@@ -175,8 +175,8 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_FUNCTION_MAP{
   {"cuMemcpy",                                             {"hipMemcpy_",                                              "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}},
   // no analogue
   // NOTE: Not equal to cudaMemcpy2D due to different signatures
-  {"cuMemcpy2D",                                           {"hipMemcpy2D_",                                            "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}},
-  {"cuMemcpy2D_v2",                                        {"hipMemcpy2D_",                                            "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}},
+  {"cuMemcpy2D",                                           {"hipMemcpyParam2D",                                        "", CONV_MEMORY, API_DRIVER}},
+  {"cuMemcpy2D_v2",                                        {"hipMemcpyParam2D",                                        "", CONV_MEMORY, API_DRIVER}},
   // no analogue
   // NOTE: Not equal to cudaMemcpy2DAsync due to different signatures
   {"cuMemcpy2DAsync",                                      {"hipMemcpy2DAsync_",                                       "", CONV_MEMORY, API_DRIVER, HIP_UNSUPPORTED}},

projects/hip/include/hip/hcc_detail/driver_types.h

@@ -80,22 +80,22 @@ typedef struct hipArray {
 }hipArray;
 
 typedef struct hip_Memcpy2D {
-    size_t height;
-    size_t widthInBytes;
-    hipArray* dstArray;
-    hipDeviceptr_t dstDevice;
-    void* dstHost;
-    hipMemoryType dstMemoryType;
-    size_t dstPitch;
-    size_t dstXInBytes;
-    size_t dstY;
-    hipArray* srcArray;
-    hipDeviceptr_t srcDevice;
-    const void* srcHost;
-    hipMemoryType srcMemoryType;
-    size_t srcPitch;
     size_t srcXInBytes;
     size_t srcY;
+    hipMemoryType srcMemoryType;
+    const void* srcHost;
+    hipDeviceptr_t srcDevice;
+    hipArray* srcArray;
+    size_t srcPitch;
+    size_t dstXInBytes;
+    size_t dstY;
+    hipMemoryType dstMemoryType;
+    void* dstHost;
+    hipDeviceptr_t dstDevice;
+    hipArray* dstArray;
+    size_t dstPitch;
+    size_t WidthInBytes;
+    size_t Height;
 } hip_Memcpy2D;
 
 

projects/hip/include/hip/hcc_detail/hip_runtime_api.h

@@ -1858,6 +1858,16 @@ hipError_t hipMalloc3DArray(hipArray** array, const struct hipChannelFormatDesc*
  */
 hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width,
                        size_t height, hipMemcpyKind kind);
+
+/**
+*  @brief Copies memory for 2D arrays.
+*  @param[in]   pCopy Parameters for the memory copy
+ *  @return      #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue,
+ * #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection
+ *
+ *  @see hipMemcpy, hipMemcpy2D, hipMemcpyToArray, hipMemcpy2DToArray, hipMemcpyFromArray,
+ * hipMemcpyToSymbol, hipMemcpyAsync
+*/
 hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy);
 
 /**

projects/hip/include/hip/nvcc_detail/hip_runtime_api.h

@@ -145,6 +145,12 @@ typedef enum hipChannelFormatKind {
 #define hipJitOptionFastCompile CU_JIT_FAST_COMPILE
 #define hipJitOptionNumOptions CU_JIT_NUM_OPTIONS
 
+// enum CUmemorytype redefines
+#define hipMemoryTypeHost CU_MEMORYTYPE_HOST
+#define hipMemoryTypeDevice CU_MEMORYTYPE_DEVICE
+#define hipMemoryTypeArray CU_MEMORYTYPE_ARRAY
+#define hipMemoryTypeUnified CU_MEMORYTYPE_UNIFIED
+
 typedef cudaEvent_t hipEvent_t;
 typedef cudaStream_t hipStream_t;
 typedef cudaIpcEventHandle_t hipIpcEventHandle_t;
@@ -162,6 +168,8 @@ typedef CUdeviceptr hipDeviceptr_t;
 typedef struct cudaArray hipArray;
 typedef struct cudaArray* hipArray_const_t;
 typedef cudaFuncAttributes hipFuncAttributes;
+typedef enum CUmemorytype hipMemoryType;
+typedef struct CUDA_MEMCPY2D hip_Memcpy2D;
 #define hipMemcpy3DParms cudaMemcpy3DParms
 #define hipArrayDefault cudaArrayDefault
 #define hipArrayLayered cudaArrayLayered
@@ -578,6 +586,10 @@ inline static hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src,
         cudaMemcpy2D(dst, dpitch, src, spitch, width, height, hipMemcpyKindToCudaMemcpyKind(kind)));
 }
 
+inline static hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) {
+  return hipCUDAErrorTohipError(cuMemcpy2D(hip_Memcpy2D));
+}
+
 inline static hipError_t hipMemcpy3D(const struct hipMemcpy3DParms *p)
 {
     return hipCUDAErrorTohipError(cudaMemcpy3D(p));

projects/hip/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp

@@ -71,8 +71,8 @@ bool runTest(int argc, char** argv) {
     copyParam.srcMemoryType = hipMemoryTypeHost;
     copyParam.srcHost = hData;
     copyParam.srcPitch = width * sizeof(float);
-    copyParam.widthInBytes = copyParam.srcPitch;
-    copyParam.height = height;
+    copyParam.WidthInBytes = copyParam.srcPitch;
+    copyParam.Height = height;
     hipMemcpyParam2D(&copyParam);
 
     textureReference* texref;

projects/hip/src/hip_memory.cpp

@@ -1715,8 +1715,8 @@ hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) {
     if (pCopy == nullptr) {
         e = hipErrorInvalidValue;
     }
-    e = ihipMemcpy2D(pCopy->dstArray->data, pCopy->widthInBytes, pCopy->srcHost, pCopy->srcPitch,
-                     pCopy->widthInBytes, pCopy->height, hipMemcpyDefault);
+    e = ihipMemcpy2D(pCopy->dstArray->data, pCopy->WidthInBytes, pCopy->srcHost, pCopy->srcPitch,
+                     pCopy->WidthInBytes, pCopy->Height, hipMemcpyDefault);
     return ihipLogStatus(e);
 }