|
else if ((handleBackend == RppBackend::RPP_HIP_BACKEND) && (executionBackend == RppBackend::RPP_HIP_BACKEND)) |
|
{ |
|
RpptDescPtr inputDesc = srcDescPtr; |
|
void *tempPtr = nullptr; |
|
if (srcDescPtr->c == 3) |
|
{ |
|
size_t elementSize = (srcDescPtr->dataType == RpptDataType::F32) ? 4 : |
|
(srcDescPtr->dataType == RpptDataType::F16) ? 2 : 1; |
|
size_t dataSize = dstDescPtr->strides.nStride * dstDescPtr->n * elementSize; |
|
|
|
CHECK_RETURN_STATUS(hipMalloc(&tempPtr, dataSize)); |
|
|
|
RpptSubpixelLayout srcSubpixelLayout = RpptSubpixelLayout::RGBtype; |
|
RppStatus errorStatus = rppt_color_to_greyscale(srcPtr, srcDescPtr, tempPtr, dstDescPtr, srcSubpixelLayout, rppHandle, RppBackend::RPP_HIP_BACKEND); |
|
if(errorStatus != RPP_SUCCESS) |
|
{ |
|
CHECK_RETURN_STATUS(hipFree(tempPtr)); |
|
return errorStatus; |
|
} |
|
inputDesc = dstDescPtr; |
|
} |
|
srcPtr = (tempPtr == nullptr) ? srcPtr : tempPtr; |
|
hipStreamSynchronize(handle.GetStream()); |
|
|
|
if ((srcDescPtr->dataType == RpptDataType::U8) && (dstDescPtr->dataType == RpptDataType::U8)) |
|
{ |
|
hip_exec_sobel_filter_tensor(static_cast<Rpp8u*>(srcPtr) + inputDesc->offsetInBytes, |
|
inputDesc, |
|
static_cast<Rpp8u*>(dstPtr) + dstDescPtr->offsetInBytes, |
|
dstDescPtr, |
|
sobelType, |
|
kernelSize, |
|
roiTensorPtrSrc, |
|
roiType, |
|
handle); |
|
} |
|
else if ((srcDescPtr->dataType == RpptDataType::F16) && (dstDescPtr->dataType == RpptDataType::F16)) |
|
{ |
|
hip_exec_sobel_filter_tensor((half*) (static_cast<Rpp8u*>(srcPtr) + inputDesc->offsetInBytes), |
|
inputDesc, |
|
(half*) (static_cast<Rpp8u*>(dstPtr) + dstDescPtr->offsetInBytes), |
|
dstDescPtr, |
|
sobelType, |
|
kernelSize, |
|
roiTensorPtrSrc, |
|
roiType, |
|
handle); |
|
} |
|
else if ((srcDescPtr->dataType == RpptDataType::F32) && (dstDescPtr->dataType == RpptDataType::F32)) |
|
{ |
|
hip_exec_sobel_filter_tensor((Rpp32f*) (static_cast<Rpp8u*>(srcPtr) + inputDesc->offsetInBytes), |
|
inputDesc, |
|
(Rpp32f*) (static_cast<Rpp8u*>(dstPtr) + dstDescPtr->offsetInBytes), |
|
dstDescPtr, |
|
sobelType, |
|
kernelSize, |
|
roiTensorPtrSrc, |
|
roiType, |
|
handle); |
|
} |
|
else if ((srcDescPtr->dataType == RpptDataType::I8) && (dstDescPtr->dataType == RpptDataType::I8)) |
|
{ |
|
hip_exec_sobel_filter_tensor(static_cast<Rpp8s*>(srcPtr) + inputDesc->offsetInBytes, |
|
inputDesc, |
|
static_cast<Rpp8s*>(dstPtr) + dstDescPtr->offsetInBytes, |
|
dstDescPtr, |
|
sobelType, |
|
kernelSize, |
|
roiTensorPtrSrc, |
|
roiType, |
|
handle); |
|
} |
|
else |
|
return RPP_ERROR_NOT_IMPLEMENTED; |
|
|
|
if (tempPtr != nullptr) |
|
CHECK_RETURN_STATUS(hipFree(tempPtr)); |
|
|
|
return RPP_SUCCESS; |
|
} |
Issue Description
The Sobel implementation contains issues with memory management related to the HIP path shown below:
rpp/src/modules/tensor/rppt_tensor_filter_augmentations.cpp
Lines 561 to 640 in 187caab
Resource Management Issues
tempPtris allocated usinghipMallocif the input tensor has 3 channels. If the branch which returnsRPP_ERROR_NOT_IMPLEMENTEDis taken, the function returns early without freeing the allocated memory.hip_exec_sobel_filter_tensormay take an RPP handle containing a stream reference, therefore launching asynchronously. Since there is no stream synchronization after this call,hipFreemay be called beforehip_exec_soble_filter_tensoris finished with its work. Leading to a use-after-free error.Status Propagation Issues
hip_exec_sobel_filter_tensorreturns anRppStatuswhich is not propagated to its caller:rppt_sobel_filter. This may hide potential error statuses whichhip_exec_soble_filter_tensorwould return. This issue also exists in the Host path for the Sobel filter implementation.