Skip to content
Permalink
Browse files

[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
  • Loading branch information...
emankov committed May 22, 2019
1 parent 673617f commit 8f059b0ee928d2d2f25c03c81d9f0dac3b1d5f62
@@ -861,7 +861,7 @@
| `cuMemAllocManaged` | |
| `cuMemAllocPitch` | |
| `cuMemcpy` | |
| `cuMemcpy2D` | |
| `cuMemcpy2D` | `hipMemcpyParam2D` |
| `cuMemcpy2DAsync` | |
| `cuMemcpy2DUnaligned` | |
| `cuMemcpy3D` | |
@@ -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}},
@@ -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;


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

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

0 comments on commit 8f059b0

Please sign in to comment.
You can’t perform that action at this time.