Skip to content

fix(rpp): Ensure RppStatus propagation through kernel dispatcher #11311

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
zacharyvincze merged 23 commits into from
Sep 24, 2026
Merged

fix(rpp): Ensure RppStatus propagation through kernel dispatcher #11311

Show file tree
Hide file tree
Changes from 4 commits
Commits
Show all changes
23 commits
Select commit Hold shift + click to select a range
cd7b356
Ensure RppStatus propagation in HIP exec calls
zacharyvincze Aug 26, 2026
78d8240
Propagate return codes in HOST calls
zacharyvincze Aug 26, 2026
cd90315
Merge remote-tracking branch 'origin/develop' into users/zacharyvincz…
zacharyvincze Aug 26, 2026
3cbfd2b
Remove unreachable return statements
zacharyvincze Aug 26, 2026
70f6169
Address review comments
zacharyvincze Aug 26, 2026
1a4ac6d
Merge remote-tracking branch 'origin/develop' into users/zacharyvincz…
zacharyvincze Aug 26, 2026
0fb8483
Remove instances of dead code
zacharyvincze Aug 26, 2026
dfe2617
Merge remote-tracking branch 'origin/develop' into users/zacharyvincz…
zacharyvincze Aug 26, 2026
9de1a5c
Add missing error code capture for tensor_add/subtract/divide/multipl…
zacharyvincze Aug 27, 2026
c855ddb
Merge remote-tracking branch 'origin/develop' into users/zacharyvincz…
zacharyvincze Aug 27, 2026
40550bd
Merge branch 'develop' into users/zacharyvincze/rpp-status-propagatio…
kiritigowda Aug 27, 2026
f17d250
Merge branch 'develop' into users/zacharyvincze/rpp-status-propagatio…
zacharyvincze Aug 28, 2026
f9f58aa
Merge branch 'develop' into users/zacharyvincze/rpp-status-propagatio…
kiritigowda Aug 31, 2026
b65a89b
Ensure errors return non-zero exit code on failure
zacharyvincze Sep 1, 2026
3b5d846
Merge remote-tracking branch 'origin/develop' into users/zacharyvincz…
zacharyvincze Sep 1, 2026
0b92331
Merge branch 'develop' into users/zacharyvincze/rpp-status-propagatio…
LakshmiKumar23 Sep 1, 2026
ff0f37f
Remove unsupported F16 in arithmetic operator tests
zacharyvincze Sep 2, 2026
ff2b5b2
Ensure error-code capture for misc HOST tests
zacharyvincze Sep 2, 2026
84b039d
Merge branch 'develop' into users/zacharyvincze/rpp-status-propagatio…
zacharyvincze Sep 2, 2026
b972b96
Ensure layout gets set even without toggle
zacharyvincze Sep 3, 2026
575f54c
Merge branch 'develop' into users/zacharyvincze/rpp-status-propagatio…
zacharyvincze Sep 3, 2026
73f905b
Merge branch 'develop' into users/zacharyvincze/rpp-status-propagatio…
zacharyvincze Sep 8, 2026
ad8c197
Merge branch 'develop' into users/zacharyvincze/rpp-status-propagatio…
zacharyvincze Sep 23, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
8 changes: 8 additions & 0 deletions projects/rpp/api/rppdefs.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -66,6 +66,14 @@ typedef halfhpp Rpp16f;
/*! \brief RPP maximum channels in audio tensor \ingroup group_rppdefs \page subpage_rppt */
#define RPPT_MAX_AUDIO_CHANNELS (16)

/*! \brief If a nested RPP call fails, return its \ref RppStatus from the enclosing function. Use
* when forwarding the result of a *_host_tensor or hip_exec_* helper. \ingroup group_rppdefs */
#define RPP_RETURN_IF_ERROR(expr) \
do { \
RppStatus _rpp_status = (expr); \
if (_rpp_status != RPP_SUCCESS) return _rpp_status; \
} while (0)

#if RPP_BACKEND_HIP
#include <hip/hip_runtime.h>
#define RPP_HOST_DEVICE __host__ __device__
Expand Down
1 change: 1 addition & 0 deletions projects/rpp/src/modules/tensor/hip/kernel/audio_tensor_add_tensor.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -59,6 +59,7 @@ RppStatus hip_exec_audio_tensor_add_tensor(Rpp32f* srcPtr1, Rpp32f* srcPtr2, Rpp
dim3(LOCAL_THREADS_X_1DIM, LOCAL_THREADS_Y_1DIM, LOCAL_THREADS_Z_1DIM), 0,
handle.GetStream(), srcPtr1, srcDescPtr->strides.nStride, dstPtr,
dstDescPtr->strides.nStride, srcLengthTensor, srcPtr2);
HIP_CHECK_LAUNCH_RETURN();

return RPP_SUCCESS;
}
1 change: 1 addition & 0 deletions projects/rpp/src/modules/tensor/hip/kernel/audio_tensor_mul_scalar.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -58,6 +58,7 @@ RppStatus hip_exec_audio_tensor_mul_scalar(Rpp32f* srcPtr, Rpp32f scalarValue,
dim3(LOCAL_THREADS_X_1DIM, LOCAL_THREADS_Y_1DIM, LOCAL_THREADS_Z_1DIM), 0,
handle.GetStream(), srcPtr, srcDescPtr->strides.nStride, dstPtr,
dstDescPtr->strides.nStride, srcLengthTensor, scalarValue);
HIP_CHECK_LAUNCH_RETURN();

return RPP_SUCCESS;
}
3 changes: 2 additions & 1 deletion projects/rpp/src/modules/tensor/hip/kernel/bitwise_and.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -161,7 +161,8 @@ RppStatus hip_exec_bitwise_and_tensor(Rpp8u* srcPtr1, Rpp8u* srcPtr2, RpptDescPt
Rpp8u* dstPtr, RpptDescPtr dstDescPtr,
RpptROIPtr roiTensorPtrSrc, RpptRoiType roiType,
rpp::Handle& handle) {
if (roiType == RpptRoiType::LTRB) hip_exec_roi_conversion_ltrb_to_xywh(roiTensorPtrSrc, handle);
if (roiType == RpptRoiType::LTRB)
RPP_RETURN_IF_ERROR(hip_exec_roi_conversion_ltrb_to_xywh(roiTensorPtrSrc, handle));

int globalThreads_x = (dstDescPtr->w + 7) >> 3;
int globalThreads_y = dstDescPtr->h;
Expand Down
3 changes: 2 additions & 1 deletion projects/rpp/src/modules/tensor/hip/kernel/bitwise_not.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -150,7 +150,8 @@ __global__ void bitwise_not_pln3_pkd3_hip_tensor(Rpp8u* srcPtr, uint3 srcStrides
RppStatus hip_exec_bitwise_not_tensor(Rpp8u* srcPtr, RpptDescPtr srcDescPtr, Rpp8u* dstPtr,
RpptDescPtr dstDescPtr, RpptROIPtr roiTensorPtrSrc,
RpptRoiType roiType, rpp::Handle& handle) {
if (roiType == RpptRoiType::LTRB) hip_exec_roi_conversion_ltrb_to_xywh(roiTensorPtrSrc, handle);
if (roiType == RpptRoiType::LTRB)
RPP_RETURN_IF_ERROR(hip_exec_roi_conversion_ltrb_to_xywh(roiTensorPtrSrc, handle));

int globalThreads_x = (dstDescPtr->w + 7) >> 3;
int globalThreads_y = dstDescPtr->h;
Expand Down
3 changes: 2 additions & 1 deletion projects/rpp/src/modules/tensor/hip/kernel/bitwise_or.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -161,7 +161,8 @@ RppStatus hip_exec_bitwise_or_tensor(Rpp8u* srcPtr1, Rpp8u* srcPtr2, RpptDescPtr
Rpp8u* dstPtr, RpptDescPtr dstDescPtr,
RpptROIPtr roiTensorPtrSrc, RpptRoiType roiType,
rpp::Handle& handle) {
if (roiType == RpptRoiType::LTRB) hip_exec_roi_conversion_ltrb_to_xywh(roiTensorPtrSrc, handle);
if (roiType == RpptRoiType::LTRB)
RPP_RETURN_IF_ERROR(hip_exec_roi_conversion_ltrb_to_xywh(roiTensorPtrSrc, handle));

int globalThreads_x = (dstDescPtr->w + 7) >> 3;
int globalThreads_y = dstDescPtr->h;
Expand Down
3 changes: 2 additions & 1 deletion projects/rpp/src/modules/tensor/hip/kernel/bitwise_xor.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -161,7 +161,8 @@ RppStatus hip_exec_bitwise_xor_tensor(Rpp8u* srcPtr1, Rpp8u* srcPtr2, RpptDescPt
Rpp8u* dstPtr, RpptDescPtr dstDescPtr,
RpptROIPtr roiTensorPtrSrc, RpptRoiType roiType,
rpp::Handle& handle) {
if (roiType == RpptRoiType::LTRB) hip_exec_roi_conversion_ltrb_to_xywh(roiTensorPtrSrc, handle);
if (roiType == RpptRoiType::LTRB)
RPP_RETURN_IF_ERROR(hip_exec_roi_conversion_ltrb_to_xywh(roiTensorPtrSrc, handle));

int globalThreads_x = (dstDescPtr->w + 7) >> 3;
int globalThreads_y = dstDescPtr->h;
Expand Down
10 changes: 8 additions & 2 deletions projects/rpp/src/modules/tensor/hip/kernel/blend.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -168,7 +168,8 @@ RppStatus hip_exec_blend_tensor(T* srcPtr1, T* srcPtr2, RpptDescPtr srcDescPtr,
RpptDescPtr dstDescPtr, Rpp32f* alphaTensor,
RpptROIPtr roiTensorPtrSrc, RpptRoiType roiType,
rpp::Handle& handle) {
if (roiType == RpptRoiType::LTRB) hip_exec_roi_conversion_ltrb_to_xywh(roiTensorPtrSrc, handle);
if (roiType == RpptRoiType::LTRB)
RPP_RETURN_IF_ERROR(hip_exec_roi_conversion_ltrb_to_xywh(roiTensorPtrSrc, handle));

int globalThreads_x = (dstDescPtr->strides.hStride + 7) >> 3;
int globalThreads_y = dstDescPtr->h;
Expand Down Expand Up @@ -241,7 +242,8 @@ template <typename T>
RppStatus hip_exec_blend_single_image(T* srcPtr1, T* srcPtr2, RpptDescPtr srcDescPtr, T* dstPtr,
RpptDescPtr dstDescPtr, Rpp32f* alphaTensor,
RpptROIPtr roiSrc, RpptRoiType roiType, rpp::Handle& handle) {
if (roiType == RpptRoiType::LTRB) hip_exec_roi_conversion_ltrb_to_xywh(roiSrc, handle);
if (roiType == RpptRoiType::LTRB)
RPP_RETURN_IF_ERROR(hip_exec_roi_conversion_ltrb_to_xywh(roiSrc, handle));

int globalThreads_x = (dstDescPtr->strides.hStride + 7) >> 3;
int globalThreads_y = dstDescPtr->h;
Expand All @@ -258,6 +260,7 @@ RppStatus hip_exec_blend_single_image(T* srcPtr1, T* srcPtr2, RpptDescPtr srcDes
srcPtr2, make_uint2(srcDescPtr->strides.nStride, srcDescPtr->strides.hStride), dstPtr,
Comment thread
zacharyvincze marked this conversation as resolved.
Resolved
make_uint2(dstDescPtr->strides.nStride, dstDescPtr->strides.hStride), alphaTensor,
roiSrc);
HIP_CHECK_LAUNCH_RETURN();
if (hipGetLastError() != hipSuccess) status = RPP_ERROR_HIP_LAUNCH;
} else if ((srcDescPtr->layout == RpptLayout::NCHW) &&
(dstDescPtr->layout == RpptLayout::NCHW)) {
Expand All @@ -273,6 +276,7 @@ RppStatus hip_exec_blend_single_image(T* srcPtr1, T* srcPtr2, RpptDescPtr srcDes
make_uint3(dstDescPtr->strides.nStride, dstDescPtr->strides.cStride,
dstDescPtr->strides.hStride),
dstDescPtr->c, alphaTensor, roiSrc);
HIP_CHECK_LAUNCH_RETURN();
if (hipGetLastError() != hipSuccess) status = RPP_ERROR_HIP_LAUNCH;
} else if ((srcDescPtr->c == 3) && (dstDescPtr->c == 3)) {
if ((srcDescPtr->layout == RpptLayout::NHWC) && (dstDescPtr->layout == RpptLayout::NCHW)) {
Expand All @@ -287,6 +291,7 @@ RppStatus hip_exec_blend_single_image(T* srcPtr1, T* srcPtr2, RpptDescPtr srcDes
make_uint3(dstDescPtr->strides.nStride, dstDescPtr->strides.cStride,
dstDescPtr->strides.hStride),
alphaTensor, roiSrc);
HIP_CHECK_LAUNCH_RETURN();
if (hipGetLastError() != hipSuccess) status = RPP_ERROR_HIP_LAUNCH;
} else if ((srcDescPtr->layout == RpptLayout::NCHW) &&
(dstDescPtr->layout == RpptLayout::NHWC)) {
Expand All @@ -302,6 +307,7 @@ RppStatus hip_exec_blend_single_image(T* srcPtr1, T* srcPtr2, RpptDescPtr srcDes
dstPtr,
make_uint2(dstDescPtr->strides.nStride, dstDescPtr->strides.hStride),
alphaTensor, roiSrc);
HIP_CHECK_LAUNCH_RETURN();
if (hipGetLastError() != hipSuccess) status = RPP_ERROR_HIP_LAUNCH;
}
}
Expand Down
6 changes: 4 additions & 2 deletions projects/rpp/src/modules/tensor/hip/kernel/box_filter.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1894,7 +1894,8 @@ RppStatus hip_exec_box_filter_tensor(T* srcPtr, RpptDescPtr srcDescPtr, T* dstPt
RpptDescPtr dstDescPtr, Rpp32u kernelSize,
RpptROIPtr roiTensorPtrSrc, RpptRoiType roiType,
rpp::Handle& handle) {
if (roiType == RpptRoiType::LTRB) hip_exec_roi_conversion_ltrb_to_xywh(roiTensorPtrSrc, handle);
if (roiType == RpptRoiType::LTRB)
RPP_RETURN_IF_ERROR(hip_exec_roi_conversion_ltrb_to_xywh(roiTensorPtrSrc, handle));

int globalThreads_x = (dstDescPtr->strides.hStride + 7) >> 3;
int globalThreads_y = dstDescPtr->h;
Expand Down Expand Up @@ -2047,7 +2048,8 @@ RppStatus hip_exec_box_filter_single_image(T* srcPtr, RpptDescPtr srcDescPtr, T*
RpptDescPtr dstDescPtr, Rpp32u kernelSize,
RpptROIPtr roiTensorPtrSrc, RpptRoiType roiType,
rpp::Handle& handle) {
if (roiType == RpptRoiType::LTRB) hip_exec_roi_conversion_ltrb_to_xywh(roiTensorPtrSrc, handle);
if (roiType == RpptRoiType::LTRB)
RPP_RETURN_IF_ERROR(hip_exec_roi_conversion_ltrb_to_xywh(roiTensorPtrSrc, handle));

int globalThreads_x = (dstDescPtr->strides.hStride + 7) >> 3;
int globalThreads_y = dstDescPtr->h;
Expand Down
6 changes: 4 additions & 2 deletions projects/rpp/src/modules/tensor/hip/kernel/brightness.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -193,7 +193,8 @@ RppStatus hip_exec_brightness_tensor(T* srcPtr, RpptDescPtr srcDescPtr, T* dstPt
RpptDescPtr dstDescPtr, Rpp32f* alphaTensor,
Rpp32f* betaTensor, RpptROIPtr roiTensorPtrSrc,
RpptRoiType roiType, rpp::Handle& handle) {
if (roiType == RpptRoiType::LTRB) hip_exec_roi_conversion_ltrb_to_xywh(roiTensorPtrSrc, handle);
if (roiType == RpptRoiType::LTRB)
RPP_RETURN_IF_ERROR(hip_exec_roi_conversion_ltrb_to_xywh(roiTensorPtrSrc, handle));

int globalThreads_x = (dstDescPtr->strides.hStride + 7) >> 3;
int globalThreads_y = dstDescPtr->h;
Expand Down Expand Up @@ -265,7 +266,8 @@ RppStatus hip_exec_brightness_single_image(T* srcPtr, RpptDescPtr srcDescPtr, T*
RpptDescPtr dstDescPtr, Rpp32f* alphaTensor,
Rpp32f* betaTensor, RpptROIPtr roiSrc,
RpptRoiType roiType, rpp::Handle& handle) {
if (roiType == RpptRoiType::LTRB) hip_exec_roi_conversion_ltrb_to_xywh(roiSrc, handle);
if (roiType == RpptRoiType::LTRB)
RPP_RETURN_IF_ERROR(hip_exec_roi_conversion_ltrb_to_xywh(roiSrc, handle));

int globalThreads_x = (dstDescPtr->strides.hStride + 7) >> 3;
int globalThreads_y = dstDescPtr->h;
Expand Down
3 changes: 2 additions & 1 deletion projects/rpp/src/modules/tensor/hip/kernel/channel_dropout.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -210,7 +210,8 @@ RppStatus hip_exec_channel_dropout_tensor(T* srcPtr, RpptDescPtr srcDescPtr, T*
RpptDescPtr dstDescPtr, Rpp8u* dropoutTensor,
RpptROIPtr roiTensorPtrSrc, RpptRoiType roiType,
rpp::Handle& handle) {
if (roiType == RpptRoiType::LTRB) hip_exec_roi_conversion_ltrb_to_xywh(roiTensorPtrSrc, handle);
if (roiType == RpptRoiType::LTRB)
RPP_RETURN_IF_ERROR(hip_exec_roi_conversion_ltrb_to_xywh(roiTensorPtrSrc, handle));

int globalThreads_x = (dstDescPtr->strides.hStride + 7) >> 3;
int globalThreads_y = dstDescPtr->h;
Expand Down
9 changes: 8 additions & 1 deletion projects/rpp/src/modules/tensor/hip/kernel/coarse_dropout.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -153,7 +153,8 @@ RppStatus hip_exec_coarse_dropout_tensor(T* srcPtr, RpptDescPtr srcDescPtr, T* d
Rpp32u* numBoxesTensor, Rpp32u maxBoxesPerImage,
RpptROIPtr roiTensorPtrSrc, RpptRoiType roiType,
rpp::Handle& handle) {
if (roiType == RpptRoiType::LTRB) hip_exec_roi_conversion_ltrb_to_xywh(roiTensorPtrSrc, handle);
if (roiType == RpptRoiType::LTRB)
RPP_RETURN_IF_ERROR(hip_exec_roi_conversion_ltrb_to_xywh(roiTensorPtrSrc, handle));

int globalThreads_x = dstDescPtr->w;
int globalThreads_y = dstDescPtr->h;
Expand Down Expand Up @@ -187,6 +188,7 @@ RppStatus hip_exec_coarse_dropout_tensor(T* srcPtr, RpptDescPtr srcDescPtr, T* d
dstPtr,
make_uint2(dstDescPtr->strides.nStride, dstDescPtr->strides.hStride),
roiTensorPtrSrc);
HIP_CHECK_LAUNCH_RETURN();
globalThreads_x = dstDescPtr->w;
RPP_HIP_RETURN_IF_ERROR(hipStreamSynchronize(handle.GetStream()));
}
Expand All @@ -199,6 +201,7 @@ RppStatus hip_exec_coarse_dropout_tensor(T* srcPtr, RpptDescPtr srcDescPtr, T* d
handle.GetStream(), dstPtr,
make_uint2(dstDescPtr->strides.nStride, dstDescPtr->strides.hStride),
anchorBoxInfoTensor, numBoxesTensor, roiTensorPtrSrc, maxBoxesPerImage);
HIP_CHECK_LAUNCH_RETURN();
} else if ((srcDescPtr->layout == RpptLayout::NCHW) &&
(dstDescPtr->layout == RpptLayout::NCHW) && dstDescPtr->c == 1) {
RPP_HIP_RETURN_IF_ERROR(hipMemcpyAsync(
Expand All @@ -215,6 +218,7 @@ RppStatus hip_exec_coarse_dropout_tensor(T* srcPtr, RpptDescPtr srcDescPtr, T* d
make_uint3(dstDescPtr->strides.nStride, dstDescPtr->strides.cStride,
dstDescPtr->strides.hStride),
anchorBoxInfoTensor, numBoxesTensor, roiTensorPtrSrc, maxBoxesPerImage);
HIP_CHECK_LAUNCH_RETURN();
} else if ((srcDescPtr->layout == RpptLayout::NCHW) &&
(dstDescPtr->layout == RpptLayout::NCHW) && dstDescPtr->c == 3) {
RPP_HIP_RETURN_IF_ERROR(hipMemcpyAsync(
Expand All @@ -231,6 +235,7 @@ RppStatus hip_exec_coarse_dropout_tensor(T* srcPtr, RpptDescPtr srcDescPtr, T* d
make_uint3(dstDescPtr->strides.nStride, dstDescPtr->strides.cStride,
dstDescPtr->strides.hStride),
anchorBoxInfoTensor, numBoxesTensor, roiTensorPtrSrc, maxBoxesPerImage);
HIP_CHECK_LAUNCH_RETURN();
} else if ((srcDescPtr->c == 3) && (dstDescPtr->c == 3)) {
if ((srcDescPtr->layout == RpptLayout::NHWC) && (dstDescPtr->layout == RpptLayout::NCHW)) {
globalThreads_x = (dstDescPtr->w + 7) >> 3;
Expand All @@ -245,6 +250,7 @@ RppStatus hip_exec_coarse_dropout_tensor(T* srcPtr, RpptDescPtr srcDescPtr, T* d
make_uint3(dstDescPtr->strides.nStride, dstDescPtr->strides.cStride,
dstDescPtr->strides.hStride),
roiTensorPtrSrc);
HIP_CHECK_LAUNCH_RETURN();
RPP_HIP_RETURN_IF_ERROR(hipStreamSynchronize(handle.GetStream()));
globalThreads_x = dstDescPtr->w;
hipLaunchKernelGGL(coarse_dropout_pln3_hip_tensor,
Expand All @@ -257,6 +263,7 @@ RppStatus hip_exec_coarse_dropout_tensor(T* srcPtr, RpptDescPtr srcDescPtr, T* d
dstDescPtr->strides.hStride),
anchorBoxInfoTensor, numBoxesTensor, roiTensorPtrSrc,
maxBoxesPerImage);
HIP_CHECK_LAUNCH_RETURN();
}
}

Expand Down
3 changes: 2 additions & 1 deletion projects/rpp/src/modules/tensor/hip/kernel/color_cast.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -183,7 +183,8 @@ RppStatus hip_exec_color_cast_tensor(T* srcPtr, RpptDescPtr srcDescPtr, T* dstPt
RpptDescPtr dstDescPtr, RpptRGB* rgbTensor,
Rpp32f* alphaTensor, RpptROIPtr roiTensorPtrSrc,
RpptRoiType roiType, rpp::Handle& handle) {
if (roiType == RpptRoiType::LTRB) hip_exec_roi_conversion_ltrb_to_xywh(roiTensorPtrSrc, handle);
if (roiType == RpptRoiType::LTRB)
RPP_RETURN_IF_ERROR(hip_exec_roi_conversion_ltrb_to_xywh(roiTensorPtrSrc, handle));

if ((srcDescPtr->c == 3) && (dstDescPtr->c == 3)) {
int globalThreads_x = (dstDescPtr->strides.hStride + 7) >> 3;
Expand Down
3 changes: 3 additions & 0 deletions projects/rpp/src/modules/tensor/hip/kernel/color_space_conversion.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -275,6 +275,7 @@ RppStatus hip_exec_yuv_to_rgb(T* srcYPtr, Rpp32u src_y_pitch, T* srcUVPtr, Rpp32
handle.GetStream(), (uint8_t*)srcYPtr, (int)src_y_pitch, (uint8_t*)srcUVPtr,
(int)src_uv_pitch, (uint8_t*)dstPtr, (int)dst_pitch, (int)width,
(int)height);
HIP_CHECK_LAUNCH_RETURN();
return RPP_SUCCESS;
}

Expand All @@ -297,6 +298,7 @@ RppStatus hip_exec_yuv_to_rgb_cubic_v(T* srcYPtr, Rpp32u src_y_pitch, T* srcUVPt
handle.GetStream(), (uint8_t*)srcYPtr, (int)src_y_pitch, (uint8_t*)srcUVPtr,
(int)src_uv_pitch, (uint8_t*)dstPtr, (int)dst_pitch, (int)width,
(int)height);
HIP_CHECK_LAUNCH_RETURN();
return RPP_SUCCESS;
}

Expand All @@ -320,6 +322,7 @@ RppStatus hip_exec_yuv_to_rgb_linear_v(T* srcYPtr, Rpp32u src_y_pitch, T* srcUVP
handle.GetStream(), (uint8_t*)srcYPtr, (int)src_y_pitch, (uint8_t*)srcUVPtr,
(int)src_uv_pitch, (uint8_t*)dstPtr, (int)dst_pitch, (int)width,
(int)height);
HIP_CHECK_LAUNCH_RETURN();
return RPP_SUCCESS;
}

Expand Down
3 changes: 2 additions & 1 deletion projects/rpp/src/modules/tensor/hip/kernel/color_temperature.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -162,7 +162,8 @@ RppStatus hip_exec_color_temperature_tensor(T* srcPtr, RpptDescPtr srcDescPtr, T
RpptDescPtr dstDescPtr, Rpp32s* adjustmentValueTensor,
RpptROIPtr roiTensorPtrSrc, RpptRoiType roiType,
rpp::Handle& handle) {
if (roiType == RpptRoiType::LTRB) hip_exec_roi_conversion_ltrb_to_xywh(roiTensorPtrSrc, handle);
if (roiType == RpptRoiType::LTRB)
RPP_RETURN_IF_ERROR(hip_exec_roi_conversion_ltrb_to_xywh(roiTensorPtrSrc, handle));

if ((srcDescPtr->c == 3) && (dstDescPtr->c == 3)) {
int globalThreads_x = (dstDescPtr->strides.hStride + 7) >> 3;
Expand Down
3 changes: 2 additions & 1 deletion projects/rpp/src/modules/tensor/hip/kernel/color_twist.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -233,7 +233,8 @@ RppStatus hip_exec_color_twist_tensor(T* srcPtr, RpptDescPtr srcDescPtr, T* dstP
Rpp32f* contrastTensor, Rpp32f* hueTensor,
Rpp32f* saturationTensor, RpptROIPtr roiTensorPtrSrc,
RpptRoiType roiType, rpp::Handle& handle) {
if (roiType == RpptRoiType::LTRB) hip_exec_roi_conversion_ltrb_to_xywh(roiTensorPtrSrc, handle);
if (roiType == RpptRoiType::LTRB)
RPP_RETURN_IF_ERROR(hip_exec_roi_conversion_ltrb_to_xywh(roiTensorPtrSrc, handle));

if ((srcDescPtr->c == 3) && (dstDescPtr->c == 3)) {
int globalThreads_x = (dstDescPtr->strides.hStride + 7) >> 3;
Expand Down
3 changes: 2 additions & 1 deletion projects/rpp/src/modules/tensor/hip/kernel/contrast.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -202,7 +202,8 @@ RppStatus hip_exec_contrast_tensor(T* srcPtr, RpptDescPtr srcDescPtr, T* dstPtr,
RpptDescPtr dstDescPtr, Rpp32f* contrastFactorTensor,
Rpp32f* contrastCenterTensor, RpptROIPtr roiTensorPtrSrc,
RpptRoiType roiType, rpp::Handle& handle) {
if (roiType == RpptRoiType::LTRB) hip_exec_roi_conversion_ltrb_to_xywh(roiTensorPtrSrc, handle);
if (roiType == RpptRoiType::LTRB)
RPP_RETURN_IF_ERROR(hip_exec_roi_conversion_ltrb_to_xywh(roiTensorPtrSrc, handle));

int globalThreads_x = (dstDescPtr->strides.hStride + 7) >> 3;
int globalThreads_y = dstDescPtr->h;
Expand Down
6 changes: 4 additions & 2 deletions projects/rpp/src/modules/tensor/hip/kernel/crop.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -126,7 +126,8 @@ template <typename T>
RppStatus hip_exec_crop_tensor(T* srcPtr, RpptDescPtr srcDescPtr, T* dstPtr, RpptDescPtr dstDescPtr,
RpptROIPtr roiTensorPtrSrc, RpptRoiType roiType,
rpp::Handle& handle) {
if (roiType == RpptRoiType::LTRB) hip_exec_roi_conversion_ltrb_to_xywh(roiTensorPtrSrc, handle);
if (roiType == RpptRoiType::LTRB)
RPP_RETURN_IF_ERROR(hip_exec_roi_conversion_ltrb_to_xywh(roiTensorPtrSrc, handle));

int globalThreads_x = (dstDescPtr->strides.hStride + 7) >> 3;
int globalThreads_y = dstDescPtr->h;
Expand Down Expand Up @@ -196,7 +197,8 @@ template <typename T>
RppStatus hip_exec_crop_single_image(T* srcPtr, RpptDescPtr srcDescPtr, T* dstPtr,
RpptDescPtr dstDescPtr, RpptROIPtr roiSrc, RpptRoiType roiType,
rpp::Handle& handle) {
if (roiType == RpptRoiType::LTRB) hip_exec_roi_conversion_ltrb_to_xywh(roiSrc, handle);
if (roiType == RpptRoiType::LTRB)
RPP_RETURN_IF_ERROR(hip_exec_roi_conversion_ltrb_to_xywh(roiSrc, handle));

int globalThreads_x = (dstDescPtr->strides.hStride + 7) >> 3;
int globalThreads_y = dstDescPtr->h;
Expand Down
Loading
Loading