From 204043c6e043515adedf979318d21d37acf3292c Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Wed, 22 May 2019 18:31:39 +0300 Subject: [PATCH 1/3] [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/clr commit: 9cb3e9aa5ef564d84d3266f1794d7771422debd0] --- ...A_Driver_API_functions_supported_by_HIP.md | 2 +- .../src/CUDA2HIP_Driver_API_functions.cpp | 4 +-- .../include/hip/hcc_detail/driver_types.h | 28 +++++++++---------- .../include/hip/hcc_detail/hip_runtime_api.h | 10 +++++++ .../include/hip/nvcc_detail/hip_runtime_api.h | 12 ++++++++ .../11_texture_driver/texture2dDrv.cpp | 4 +-- projects/clr/hipamd/src/hip_memory.cpp | 4 +-- 7 files changed, 43 insertions(+), 21 deletions(-) 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..08897fdb07 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 @@ -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)); 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); } From a0e1887ff3f63b1896efbe9b3f7e1fb3989dacf8 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Wed, 22 May 2019 20:04:45 +0300 Subject: [PATCH 2/3] [HIP] fix nvcc path break in #1127 [ROCm/clr commit: 6806ab6745374d83a6f630f84fcc2e9d0d88c580] --- .../hipamd/include/hip/nvcc_detail/hip_runtime_api.h | 11 ++--------- 1 file changed, 2 insertions(+), 9 deletions(-) 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 08897fdb07..7bc7b91368 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 @@ -145,12 +145,6 @@ 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; @@ -168,8 +162,7 @@ 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 hip_Memcpy2D CUDA_MEMCPY2D #define hipMemcpy3DParms cudaMemcpy3DParms #define hipArrayDefault cudaArrayDefault #define hipArrayLayered cudaArrayLayered @@ -587,7 +580,7 @@ inline static hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, } inline static hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) { - return hipCUDAErrorTohipError(cuMemcpy2D(hip_Memcpy2D)); + return hipCUResultTohipError(cuMemcpy2D(hip_Memcpy2D)); } inline static hipError_t hipMemcpy3D(const struct hipMemcpy3DParms *p) From 3afaf0d2de957d02a429363b7cb236a016264363 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Wed, 22 May 2019 20:48:18 +0300 Subject: [PATCH 3/3] [HIP] fix typo in #1127 [ROCm/clr commit: 49b9df7a9e78dc72ad97148ce1b3aa20c569490f] --- projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) 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 7bc7b91368..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 @@ -580,7 +580,7 @@ inline static hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, } inline static hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) { - return hipCUResultTohipError(cuMemcpy2D(hip_Memcpy2D)); + return hipCUResultTohipError(cuMemcpy2D(pCopy)); } inline static hipError_t hipMemcpy3D(const struct hipMemcpy3DParms *p)