From 2ca3e8f3371056f88a4623c1994dafe617ce2e65 Mon Sep 17 00:00:00 2001 From: foreman Date: Wed, 25 Sep 2019 16:53:50 -0400 Subject: [PATCH] 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: e37b6b6740c6e769f80db634d1601f18ebca32b1] --- projects/hip/api/hip/hip_memory.cpp | 24 +++++++++++++++--------- 1 file changed, 15 insertions(+), 9 deletions(-) diff --git a/projects/hip/api/hip/hip_memory.cpp b/projects/hip/api/hip/hip_memory.cpp index 00b75af5b1..19262ecc54 100644 --- a/projects/hip/api/hip/hip_memory.cpp +++ b/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);