diff --git a/projects/clr/hipamd/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md b/projects/clr/hipamd/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md index aeb834c2ec..7c2d2c6631 100644 --- a/projects/clr/hipamd/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md +++ b/projects/clr/hipamd/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md @@ -861,7 +861,7 @@ | `cuMemAllocManaged` | | | `cuMemAllocPitch` | | | `cuMemcpy` | | -| `cuMemcpy2D` | | +| `cuMemcpy2D` | `hipMemcpyParam2D` | | `cuMemcpy2DAsync` | | | `cuMemcpy2DUnaligned` | | | `cuMemcpy3D` | | diff --git a/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp b/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp index b71e49710d..8be20774ea 100644 --- a/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp +++ b/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Driver_API_functions.cpp @@ -175,8 +175,8 @@ const std::map 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}}, diff --git a/projects/clr/hipamd/include/hip/hcc_detail/driver_types.h b/projects/clr/hipamd/include/hip/hcc_detail/driver_types.h index 8e1fec11fa..5b2297114f 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/driver_types.h +++ b/projects/clr/hipamd/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; diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h index 5b598b54a8..d870963101 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/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); /** diff --git a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h index 2ebd4f8b7d..c1846c1b1e 100644 --- a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h @@ -162,6 +162,7 @@ typedef CUdeviceptr hipDeviceptr_t; typedef struct cudaArray hipArray; typedef struct cudaArray* hipArray_const_t; typedef cudaFuncAttributes hipFuncAttributes; +#define hip_Memcpy2D CUDA_MEMCPY2D #define hipMemcpy3DParms cudaMemcpy3DParms #define hipArrayDefault cudaArrayDefault #define hipArrayLayered cudaArrayLayered @@ -578,6 +579,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 hipCUResultTohipError(cuMemcpy2D(pCopy)); +} + inline static hipError_t hipMemcpy3D(const struct hipMemcpy3DParms *p) { return hipCUDAErrorTohipError(cudaMemcpy3D(p)); diff --git a/projects/clr/hipamd/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp b/projects/clr/hipamd/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp index a6ba44696a..36d37a4fad 100644 --- a/projects/clr/hipamd/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp +++ b/projects/clr/hipamd/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(©Param); textureReference* texref; diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index 0fad8ab890..36edcdb338 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/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); }