P4 to Git Change 2004245 by axie@axie_win_opencl_nvme on 2019/09/25 16:46:31

SWDEV-203855 - Segfault when using hipArrayCreate and hipMemcpyParam2D
	1. hipArrayCreate API implementation uses a wrong parameter to check width. That parameter can be null pointer because it is used to pass the pointer back to the caller.
	2. Implement hipMemcpyParam2D similar to HIP-HCC implementation. Reference: https://github.com/ROCm-Developer-Tools/HIP/blob/master/src/hip_memory.cpp

	Tests:
	1. PRE CHECK-IN build and test(no regression):  http://ocltc:8111/viewModification.html?modId=126608&personal=true&init=1&tab=vcsModificationBuilds
	2. GPU is VEGA10, OS is Windows 10, CPU is threadripper 1900x, run the test. There is not segfault or exit during hipArrayCreate and hipMemcpyParam2D function call.

	ReviewBoard: http://ocltc.amd.com/reviews/r/18037/

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#72 edit


[ROCm/hip commit: e37b6b6740]

projects/hip/api/hip/hip_memory.cpp

@@ -429,7 +429,7 @@ hipError_t hipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent)
 hipError_t hipArrayCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocateArray) {
   HIP_INIT_API(array, pAllocateArray);
 
-  if (array[0]->width == 0) {
+  if (pAllocateArray->Width == 0) {
     HIP_RETURN(hipErrorInvalidValue);
   }
 
@@ -815,14 +815,6 @@ hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes,
                    *queue, true));
 }
 
-hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) {
-  HIP_INIT_API(pCopy);
-
-  assert(0 && "Unimplemented");
-
-  HIP_RETURN(hipErrorUnknown);
-}
-
 hipError_t ihipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width,
                         size_t height, hipMemcpyKind kind, amd::HostQueue& queue,
                         bool isAsync = false) {
@@ -889,6 +881,20 @@ hipError_t ihipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch
   return hipSuccess;
 }
 
+hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) {
+  HIP_INIT_API(pCopy);
+  hipError_t e = hipSuccess;
+  if (pCopy == nullptr) {
+    e = hipErrorInvalidValue;
+  } else {
+    hip::syncStreams();
+    amd::HostQueue* queue = hip::getNullStream();
+    e = ihipMemcpy2D(pCopy->dstArray->data, pCopy->WidthInBytes, pCopy->srcHost, pCopy->srcPitch,
+                     pCopy->WidthInBytes, pCopy->Height, hipMemcpyDefault, *queue);
+  }
+  HIP_RETURN(e);
+}
+
 hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width,
                        size_t height, hipMemcpyKind kind) {
   HIP_INIT_API(dst, dpitch, src, spitch, width, height, kind);