From 8a52bf1051103a3927e55f70c6fb9e04a309aa37 Mon Sep 17 00:00:00 2001 From: LokeshBonta Date: Thu, 28 May 2020 04:25:34 +0530 Subject: [PATCH] Support for ColorTwist in Float space --- src/modules/cl/cl_fused_functions.cpp | 2 +- src/modules/cl/kernel/colortwist.cl | 36 ++- src/modules/rppi_fused_functions.cpp | 425 +++++--------------------- 3 files changed, 89 insertions(+), 374 deletions(-) diff --git a/src/modules/cl/cl_fused_functions.cpp b/src/modules/cl/cl_fused_functions.cpp index 3fe2e670b..59e8df87e 100644 --- a/src/modules/cl/cl_fused_functions.cpp +++ b/src/modules/cl/cl_fused_functions.cpp @@ -39,7 +39,7 @@ color_twist_cl_batch ( cl_mem srcPtr, cl_mem dstPtr, rpp::Handle& handle, std::vector vld{16, 16, 1}; std::vector vgd{max_width , max_height, handle.GetBatchSize()}; std::string kernel_file = "colortwist.cl"; - std::string kernel_name = "colortwisth_batch"; + std::string kernel_name = "colortwist_batch"; switch (dataType) { case RPPTensorDataType::U8: diff --git a/src/modules/cl/kernel/colortwist.cl b/src/modules/cl/kernel/colortwist.cl index 3fbbe386c..7fe63174e 100644 --- a/src/modules/cl/kernel/colortwist.cl +++ b/src/modules/cl/kernel/colortwist.cl @@ -159,8 +159,6 @@ __kernel void colortwist_pln( __global unsigned char *input, output[pixIdx + 2*width*height] = saturate_8u(alpha * pixel.z + beta); } - - __kernel void colortwist_batch( __global unsigned char* input, __global unsigned char* output, __global float *alpha, @@ -212,7 +210,7 @@ __kernel void colortwist_batch( __global unsigned char* input, } } -__kernel void colortwist_batch_fp32(__global float* input, +__kernel void colortwist_batch_fp32( __global float* input, __global float* output, __global float *alpha, __global float *beta, @@ -232,13 +230,13 @@ __kernel void colortwist_batch_fp32(__global float* input, { int id_x = get_global_id(0), id_y = get_global_id(1), id_z = get_global_id(2); if (id_x >= width[id_z] || id_y >= height[id_z]) return; - float4 pixel; float4 hsv; + uchar4 pixel; float4 hsv; unsigned int l_inc = inc[id_z]; // for local increment int pixIdx = batch_index[id_z] + (id_y * max_width[id_z] + id_x) * plnpkdindex; - pixel.x = input[pixIdx]; - pixel.y = input[pixIdx + l_inc]; - pixel.z = input[pixIdx + 2*l_inc]; + pixel.x = (uchar)(input[pixIdx] * 255); + pixel.y = (uchar)(input[pixIdx + l_inc] * 255); + pixel.z = (uchar)(input[pixIdx + 2*l_inc] * 255); pixel.w = 0.0; float alpha1 = alpha[id_z], beta1 = beta[id_z]; @@ -257,13 +255,13 @@ __kernel void colortwist_batch_fp32(__global float* input, output[pixIdx + 2*l_inc] = saturate_8u(alpha1 * pixel.z + beta1); } else { - output[pixIdx] = pixel.x; - output[pixIdx + l_inc] = pixel.y; - output[pixIdx + 2*l_inc] = pixel.z; + output[pixIdx] = (float) (pixel.x / 255.0); + output[pixIdx + l_inc] = (float) (pixel.y / 255.0); + output[pixIdx + 2*l_inc] = (float) (pixel.z / 255.0); } - } -__kernel void colortwist_batch_fp16(__global half* input, + +__kernel void colortwist_batch_fp16( __global half* input, __global half* output, __global float *alpha, __global float *beta, @@ -287,9 +285,9 @@ __kernel void colortwist_batch_fp16(__global half* input, unsigned int l_inc = inc[id_z]; // for local increment int pixIdx = batch_index[id_z] + (id_y * max_width[id_z] + id_x) * plnpkdindex; - pixel.x = input[pixIdx]; - pixel.y = input[pixIdx + l_inc]; - pixel.z = input[pixIdx + 2*l_inc]; + pixel.x = (uchar)(input[pixIdx] * 255); + pixel.y = (uchar)(input[pixIdx + l_inc] * 255); + pixel.z = (uchar)(input[pixIdx + 2*l_inc] * 255); pixel.w = 0.0; float alpha1 = alpha[id_z], beta1 = beta[id_z]; @@ -308,10 +306,10 @@ __kernel void colortwist_batch_fp16(__global half* input, output[pixIdx + 2*l_inc] = saturate_8u(alpha1 * pixel.z + beta1); } else { - output[pixIdx] = pixel.x; - output[pixIdx + l_inc] = pixel.y; - output[pixIdx + 2*l_inc] = pixel.z; + output[pixIdx] = (half) (pixel.x / 255.0); + output[pixIdx + l_inc] = (half) (pixel.y / 255.0); + output[pixIdx + 2*l_inc] = (half) (pixel.z / 255.0); } - } + diff --git a/src/modules/rppi_fused_functions.cpp b/src/modules/rppi_fused_functions.cpp index 48ead89f8..7828a8f9a 100644 --- a/src/modules/rppi_fused_functions.cpp +++ b/src/modules/rppi_fused_functions.cpp @@ -20,6 +20,50 @@ using namespace std::chrono; #include "cpu/host_fused_functions.hpp" +RppStatus color_twist_helper(RppiChnFormat chn_format, Rpp32u num_of_channels, + RPPTensorDataType tensor_type, + RppPtr_t srcPtr ,RppiSize *srcSize ,RppiSize maxSrcSize ,RppPtr_t dstPtr ,Rpp32f *alpha ,Rpp32f *beta , + Rpp32f *hueShift ,Rpp32f *saturationFactor ,Rpp32u nbatchSize ,rppHandle_t rppHandle ) +{ + make_data_type(tensor_type); + RppiROI roiPoints; + roiPoints.x = 0; + roiPoints.y = 0; + roiPoints.roiHeight = 0; + roiPoints.roiWidth = 0; + Rpp32u paramIndex = 0; + copy_srcSize(srcSize, rpp::deref(rppHandle)); + copy_srcMaxSize (maxSrcSize, rpp::deref(rppHandle)); + copy_roi(roiPoints, rpp::deref(rppHandle)); + get_srcBatchIndex (rpp::deref(rppHandle), num_of_channels, chn_format); + copy_param_float (alpha, rpp::deref(rppHandle), paramIndex++); + copy_param_float (beta, rpp::deref(rppHandle), paramIndex++); + copy_param_float (hueShift, rpp::deref(rppHandle), paramIndex++); + copy_param_float (saturationFactor, rpp::deref(rppHandle), paramIndex++); + + #ifdef OCL_COMPILE + { + color_twist_cl_batch( + static_cast(srcPtr), + static_cast(dstPtr), + rpp::deref(rppHandle), + chn_format, num_of_channels, tensor_type + ); + } + #elif defined (HIP_COMPILE) + { + color_twist_hip_batch( + static_cast(srcPtr), + static_cast(dstPtr), + rpp::deref(rppHandle), + RPPI_CHN_PLANAR, 1 + ); + } + #endif //BACKEND + + return RPP_SUCCESS; +} + RppStatus rppi_color_twist_u8_pln1_gpu(RppPtr_t srcPtr ,RppiSize srcSize ,RppPtr_t dstPtr ,Rpp32f alpha ,Rpp32f beta ,Rpp32f hueShift ,Rpp32f saturationFactor ,rppHandle_t rppHandle ) { @@ -299,42 +343,9 @@ rppi_color_twist_u8_pln1_batchDD_gpu(RppPtr_t srcPtr ,RppiSize *srcSize ,RppPtr_ RppStatus rppi_color_twist_u8_pln1_batchPD_gpu(RppPtr_t srcPtr ,RppiSize *srcSize ,RppiSize maxSrcSize ,RppPtr_t dstPtr ,Rpp32f *alpha ,Rpp32f *beta ,Rpp32f *hueShift ,Rpp32f *saturationFactor ,Rpp32u nbatchSize ,rppHandle_t rppHandle ) { - RppiROI roiPoints; - roiPoints.x = 0; - roiPoints.y = 0; - roiPoints.roiHeight = 0; - roiPoints.roiWidth = 0; - Rpp32u paramIndex = 0; - copy_srcSize(srcSize, rpp::deref(rppHandle)); - copy_srcMaxSize (maxSrcSize, rpp::deref(rppHandle)); - copy_roi(roiPoints, rpp::deref(rppHandle)); - get_srcBatchIndex (rpp::deref(rppHandle), 1, RPPI_CHN_PLANAR); - copy_param_float (alpha, rpp::deref(rppHandle), paramIndex++); - copy_param_float (beta, rpp::deref(rppHandle), paramIndex++); - copy_param_float (hueShift, rpp::deref(rppHandle), paramIndex++); - copy_param_float (saturationFactor, rpp::deref(rppHandle), paramIndex++); - -#ifdef OCL_COMPILE - { - color_twist_cl_batch( - static_cast(srcPtr), - static_cast(dstPtr), - rpp::deref(rppHandle), - RPPI_CHN_PLANAR, 1 - ); - } -#elif defined (HIP_COMPILE) - { - color_twist_hip_batch( - static_cast(srcPtr), - static_cast(dstPtr), - rpp::deref(rppHandle), - RPPI_CHN_PLANAR, 1 - ); - } -#endif //BACKEND - - return RPP_SUCCESS; + return(color_twist_helper(RPPI_CHN_PLANAR, 1, RPPTensorDataType::U8, + srcPtr, srcSize, maxSrcSize, dstPtr, alpha, beta, + hueShift, saturationFactor, nbatchSize, rppHandle)); } RppStatus @@ -1048,42 +1059,9 @@ rppi_color_twist_u8_pln3_batchDD_gpu(RppPtr_t srcPtr ,RppiSize *srcSize ,RppPtr_ RppStatus rppi_color_twist_u8_pln3_batchPD_gpu(RppPtr_t srcPtr ,RppiSize *srcSize ,RppiSize maxSrcSize ,RppPtr_t dstPtr ,Rpp32f *alpha ,Rpp32f *beta ,Rpp32f *hueShift ,Rpp32f *saturationFactor ,Rpp32u nbatchSize ,rppHandle_t rppHandle ) { - RppiROI roiPoints; - roiPoints.x = 0; - roiPoints.y = 0; - roiPoints.roiHeight = 0; - roiPoints.roiWidth = 0; - Rpp32u paramIndex = 0; - copy_srcSize(srcSize, rpp::deref(rppHandle)); - copy_srcMaxSize (maxSrcSize, rpp::deref(rppHandle)); - copy_roi(roiPoints, rpp::deref(rppHandle)); - get_srcBatchIndex (rpp::deref(rppHandle), 3, RPPI_CHN_PLANAR); - copy_param_float (alpha, rpp::deref(rppHandle), paramIndex++); - copy_param_float (beta, rpp::deref(rppHandle), paramIndex++); - copy_param_float (hueShift, rpp::deref(rppHandle), paramIndex++); - copy_param_float (saturationFactor, rpp::deref(rppHandle), paramIndex++); - -#ifdef OCL_COMPILE - { - color_twist_cl_batch( - static_cast(srcPtr), - static_cast(dstPtr), - rpp::deref(rppHandle), - RPPI_CHN_PLANAR, 3 - ); - } -#elif defined (HIP_COMPILE) - { - color_twist_hip_batch( - static_cast(srcPtr), - static_cast(dstPtr), - rpp::deref(rppHandle), - RPPI_CHN_PLANAR, 3 - ); - } -#endif //BACKEND - - return RPP_SUCCESS; + return(color_twist_helper(RPPI_CHN_PLANAR, 3, RPPTensorDataType::U8, + srcPtr, srcSize, maxSrcSize, dstPtr, alpha, beta, + hueShift, saturationFactor, nbatchSize, rppHandle)); } RppStatus @@ -1797,42 +1775,9 @@ rppi_color_twist_u8_pkd3_batchDD_gpu(RppPtr_t srcPtr ,RppiSize *srcSize ,RppPtr_ RppStatus rppi_color_twist_u8_pkd3_batchPD_gpu(RppPtr_t srcPtr ,RppiSize *srcSize ,RppiSize maxSrcSize ,RppPtr_t dstPtr ,Rpp32f *alpha ,Rpp32f *beta ,Rpp32f *hueShift ,Rpp32f *saturationFactor ,Rpp32u nbatchSize ,rppHandle_t rppHandle ) { - RppiROI roiPoints; - roiPoints.x = 0; - roiPoints.y = 0; - roiPoints.roiHeight = 0; - roiPoints.roiWidth = 0; - Rpp32u paramIndex = 0; - copy_srcSize(srcSize, rpp::deref(rppHandle)); - copy_srcMaxSize (maxSrcSize, rpp::deref(rppHandle)); - copy_roi(roiPoints, rpp::deref(rppHandle)); - get_srcBatchIndex (rpp::deref(rppHandle), 3, RPPI_CHN_PACKED); - copy_param_float (alpha, rpp::deref(rppHandle), paramIndex++); - copy_param_float (beta, rpp::deref(rppHandle), paramIndex++); - copy_param_float (hueShift, rpp::deref(rppHandle), paramIndex++); - copy_param_float (saturationFactor, rpp::deref(rppHandle), paramIndex++); - -#ifdef OCL_COMPILE - { - color_twist_cl_batch( - static_cast(srcPtr), - static_cast(dstPtr), - rpp::deref(rppHandle), - RPPI_CHN_PACKED, 3 - ); - } -#elif defined (HIP_COMPILE) - { - color_twist_hip_batch( - static_cast(srcPtr), - static_cast(dstPtr), - rpp::deref(rppHandle), - RPPI_CHN_PACKED, 3 - ); - } -#endif //BACKEND - - return RPP_SUCCESS; + return(color_twist_helper(RPPI_CHN_PACKED, 3, RPPTensorDataType::U8, + srcPtr, srcSize, maxSrcSize, dstPtr, alpha, beta, + hueShift, saturationFactor, nbatchSize, rppHandle)); } RppStatus @@ -3921,277 +3866,49 @@ rppi_color_twist_f32_pkd3_batchPD_host(RppPtr_t srcPtr ,RppiSize *srcSize ,RppiS RppStatus rppi_color_twist_f32_pln1_batchPD_gpu(RppPtr_t srcPtr ,RppiSize *srcSize ,RppiSize maxSrcSize ,RppPtr_t dstPtr ,Rpp32f *alpha ,Rpp32f *beta ,Rpp32f *hueShift ,Rpp32f *saturationFactor ,Rpp32u nbatchSize ,rppHandle_t rppHandle ) { - RppiChnFormat chn_format = RPPI_CHN_PLANAR; - Rpp32u num_of_channels = 1; - typedef Rpp32f data_type_t; - RPPTensorDataType tensor_type = RPPTensorDataType::FP32; - - RppiROI roiPoints; - roiPoints.x = 0; - roiPoints.y = 0; - roiPoints.roiHeight = 0; - roiPoints.roiWidth = 0; - Rpp32u paramIndex = 0; - copy_srcSize(srcSize, rpp::deref(rppHandle)); - copy_srcMaxSize (maxSrcSize, rpp::deref(rppHandle)); - copy_roi(roiPoints, rpp::deref(rppHandle)); - get_srcBatchIndex (rpp::deref(rppHandle), num_of_channels, chn_format); - copy_param_float (alpha, rpp::deref(rppHandle), paramIndex++); - copy_param_float (beta, rpp::deref(rppHandle), paramIndex++); - copy_param_float (hueShift, rpp::deref(rppHandle), paramIndex++); - copy_param_float (saturationFactor, rpp::deref(rppHandle), paramIndex++); - -#ifdef OCL_COMPILE - { - color_twist_cl_batch( - static_cast(srcPtr), - static_cast(dstPtr), - rpp::deref(rppHandle), - chn_format, num_of_channels, tensor_type - ); - } -#elif defined (HIP_COMPILE) - { - color_twist_hip_batch( - static_cast(srcPtr), - static_cast(dstPtr), - rpp::deref(rppHandle), - RPPI_CHN_PLANAR, 1 - ); - } -#endif //BACKEND - - return RPP_SUCCESS; + return(color_twist_helper(RPPI_CHN_PLANAR, 1, RPPTensorDataType::FP32, + srcPtr, srcSize, maxSrcSize, dstPtr, alpha, beta, + hueShift, saturationFactor, nbatchSize, rppHandle)); } RppStatus rppi_color_twist_f32_pln3_batchPD_gpu(RppPtr_t srcPtr ,RppiSize *srcSize ,RppiSize maxSrcSize ,RppPtr_t dstPtr ,Rpp32f *alpha ,Rpp32f *beta ,Rpp32f *hueShift ,Rpp32f *saturationFactor ,Rpp32u nbatchSize ,rppHandle_t rppHandle ) { - RppiChnFormat chn_format = RPPI_CHN_PLANAR; - Rpp32u num_of_channels = 3; - typedef Rpp32f data_type_t; - RPPTensorDataType tensor_type = RPPTensorDataType::FP32; - - RppiROI roiPoints; - roiPoints.x = 0; - roiPoints.y = 0; - roiPoints.roiHeight = 0; - roiPoints.roiWidth = 0; - Rpp32u paramIndex = 0; - copy_srcSize(srcSize, rpp::deref(rppHandle)); - copy_srcMaxSize (maxSrcSize, rpp::deref(rppHandle)); - copy_roi(roiPoints, rpp::deref(rppHandle)); - get_srcBatchIndex (rpp::deref(rppHandle), num_of_channels, chn_format); - copy_param_float (alpha, rpp::deref(rppHandle), paramIndex++); - copy_param_float (beta, rpp::deref(rppHandle), paramIndex++); - copy_param_float (hueShift, rpp::deref(rppHandle), paramIndex++); - copy_param_float (saturationFactor, rpp::deref(rppHandle), paramIndex++); - -#ifdef OCL_COMPILE - { - color_twist_cl_batch( - static_cast(srcPtr), - static_cast(dstPtr), - rpp::deref(rppHandle), - chn_format, num_of_channels, tensor_type - ); - } -#elif defined (HIP_COMPILE) - { - color_twist_hip_batch( - static_cast(srcPtr), - static_cast(dstPtr), - rpp::deref(rppHandle), - chn_format, num_of_channels - ); - } -#endif //BACKEND - - return RPP_SUCCESS; + return(color_twist_helper(RPPI_CHN_PLANAR, 3, RPPTensorDataType::FP32, + srcPtr, srcSize, maxSrcSize, dstPtr, alpha, beta, + hueShift, saturationFactor, nbatchSize, rppHandle)); } RppStatus rppi_color_twist_f32_pkd3_batchPD_gpu(RppPtr_t srcPtr ,RppiSize *srcSize ,RppiSize maxSrcSize ,RppPtr_t dstPtr ,Rpp32f *alpha ,Rpp32f *beta ,Rpp32f *hueShift ,Rpp32f *saturationFactor ,Rpp32u nbatchSize ,rppHandle_t rppHandle ) { - RppiChnFormat chn_format = RPPI_CHN_PACKED; - Rpp32u num_of_channels = 3; - typedef Rpp32f data_type_t; - RPPTensorDataType tensor_type = RPPTensorDataType::FP32; - - RppiROI roiPoints; - roiPoints.x = 0; - roiPoints.y = 0; - roiPoints.roiHeight = 0; - roiPoints.roiWidth = 0; - Rpp32u paramIndex = 0; - copy_srcSize(srcSize, rpp::deref(rppHandle)); - copy_srcMaxSize (maxSrcSize, rpp::deref(rppHandle)); - copy_roi(roiPoints, rpp::deref(rppHandle)); - get_srcBatchIndex (rpp::deref(rppHandle), num_of_channels, chn_format); - copy_param_float (alpha, rpp::deref(rppHandle), paramIndex++); - copy_param_float (beta, rpp::deref(rppHandle), paramIndex++); - copy_param_float (hueShift, rpp::deref(rppHandle), paramIndex++); - copy_param_float (saturationFactor, rpp::deref(rppHandle), paramIndex++); - -#ifdef OCL_COMPILE - { - color_twist_cl_batch( - static_cast(srcPtr), - static_cast(dstPtr), - rpp::deref(rppHandle), - chn_format, num_of_channels, tensor_type - ); - } -#elif defined (HIP_COMPILE) - { - color_twist_hip_batch( - static_cast(srcPtr), - static_cast(dstPtr), - rpp::deref(rppHandle), - chn_format, num_of_channels - ); - } -#endif //BACKEND - - return RPP_SUCCESS; + return(color_twist_helper(RPPI_CHN_PACKED, 3, RPPTensorDataType::FP32, + srcPtr, srcSize, maxSrcSize, dstPtr, alpha, beta, + hueShift, saturationFactor, nbatchSize, rppHandle)); } RppStatus rppi_color_twist_f16_pln1_batchPD_gpu(RppPtr_t srcPtr ,RppiSize *srcSize ,RppiSize maxSrcSize ,RppPtr_t dstPtr ,Rpp32f *alpha ,Rpp32f *beta ,Rpp32f *hueShift ,Rpp32f *saturationFactor ,Rpp32u nbatchSize ,rppHandle_t rppHandle ) { - RppiChnFormat chn_format = RPPI_CHN_PLANAR; - Rpp32u num_of_channels = 1; - typedef Rpp16f data_type_t; - RPPTensorDataType tensor_type = RPPTensorDataType::FP16; - - RppiROI roiPoints; - roiPoints.x = 0; - roiPoints.y = 0; - roiPoints.roiHeight = 0; - roiPoints.roiWidth = 0; - Rpp32u paramIndex = 0; - copy_srcSize(srcSize, rpp::deref(rppHandle)); - copy_srcMaxSize (maxSrcSize, rpp::deref(rppHandle)); - copy_roi(roiPoints, rpp::deref(rppHandle)); - get_srcBatchIndex (rpp::deref(rppHandle), num_of_channels, chn_format); - copy_param_float (alpha, rpp::deref(rppHandle), paramIndex++); - copy_param_float (beta, rpp::deref(rppHandle), paramIndex++); - copy_param_float (hueShift, rpp::deref(rppHandle), paramIndex++); - copy_param_float (saturationFactor, rpp::deref(rppHandle), paramIndex++); - -#ifdef OCL_COMPILE - { - color_twist_cl_batch( - static_cast(srcPtr), - static_cast(dstPtr), - rpp::deref(rppHandle), - chn_format, num_of_channels, tensor_type - ); - } -#elif defined (HIP_COMPILE) - { - color_twist_hip_batch( - static_cast(srcPtr), - static_cast(dstPtr), - rpp::deref(rppHandle), - RPPI_CHN_PLANAR, 1 - ); - } -#endif //BACKEND - - return RPP_SUCCESS; + return(color_twist_helper(RPPI_CHN_PLANAR, 1, RPPTensorDataType::FP16, + srcPtr, srcSize, maxSrcSize, dstPtr, alpha, beta, + hueShift, saturationFactor, nbatchSize, rppHandle)); } RppStatus rppi_color_twist_f16_pln3_batchPD_gpu(RppPtr_t srcPtr ,RppiSize *srcSize ,RppiSize maxSrcSize ,RppPtr_t dstPtr ,Rpp32f *alpha ,Rpp32f *beta ,Rpp32f *hueShift ,Rpp32f *saturationFactor ,Rpp32u nbatchSize ,rppHandle_t rppHandle ) { - RppiChnFormat chn_format = RPPI_CHN_PLANAR; - Rpp32u num_of_channels = 3; - typedef Rpp16f data_type_t; - RPPTensorDataType tensor_type = RPPTensorDataType::FP16; - - RppiROI roiPoints; - roiPoints.x = 0; - roiPoints.y = 0; - roiPoints.roiHeight = 0; - roiPoints.roiWidth = 0; - Rpp32u paramIndex = 0; - copy_srcSize(srcSize, rpp::deref(rppHandle)); - copy_srcMaxSize (maxSrcSize, rpp::deref(rppHandle)); - copy_roi(roiPoints, rpp::deref(rppHandle)); - get_srcBatchIndex (rpp::deref(rppHandle), num_of_channels, chn_format); - copy_param_float (alpha, rpp::deref(rppHandle), paramIndex++); - copy_param_float (beta, rpp::deref(rppHandle), paramIndex++); - copy_param_float (hueShift, rpp::deref(rppHandle), paramIndex++); - copy_param_float (saturationFactor, rpp::deref(rppHandle), paramIndex++); - -#ifdef OCL_COMPILE - { - color_twist_cl_batch( - static_cast(srcPtr), - static_cast(dstPtr), - rpp::deref(rppHandle), - chn_format, num_of_channels, tensor_type - ); - } -#elif defined (HIP_COMPILE) - { - color_twist_hip_batch( - static_cast(srcPtr), - static_cast(dstPtr), - rpp::deref(rppHandle), - chn_format, num_of_channels - ); - } -#endif //BACKEND - - return RPP_SUCCESS; + return(color_twist_helper(RPPI_CHN_PLANAR, 3, RPPTensorDataType::FP16, + srcPtr, srcSize, maxSrcSize, dstPtr, alpha, beta, + hueShift, saturationFactor, nbatchSize, rppHandle)); } RppStatus rppi_color_twist_f16_pkd3_batchPD_gpu(RppPtr_t srcPtr ,RppiSize *srcSize ,RppiSize maxSrcSize ,RppPtr_t dstPtr ,Rpp32f *alpha ,Rpp32f *beta ,Rpp32f *hueShift ,Rpp32f *saturationFactor ,Rpp32u nbatchSize ,rppHandle_t rppHandle ) { - RppiChnFormat chn_format = RPPI_CHN_PACKED; - Rpp32u num_of_channels = 3; - typedef Rpp16f data_type_t; - RPPTensorDataType tensor_type = RPPTensorDataType::FP16; - - RppiROI roiPoints; - roiPoints.x = 0; - roiPoints.y = 0; - roiPoints.roiHeight = 0; - roiPoints.roiWidth = 0; - Rpp32u paramIndex = 0; - copy_srcSize(srcSize, rpp::deref(rppHandle)); - copy_srcMaxSize (maxSrcSize, rpp::deref(rppHandle)); - copy_roi(roiPoints, rpp::deref(rppHandle)); - get_srcBatchIndex (rpp::deref(rppHandle), num_of_channels, chn_format); - copy_param_float (alpha, rpp::deref(rppHandle), paramIndex++); - copy_param_float (beta, rpp::deref(rppHandle), paramIndex++); - copy_param_float (hueShift, rpp::deref(rppHandle), paramIndex++); - copy_param_float (saturationFactor, rpp::deref(rppHandle), paramIndex++); - -#ifdef OCL_COMPILE - { - color_twist_cl_batch( - static_cast(srcPtr), - static_cast(dstPtr), - rpp::deref(rppHandle), - chn_format, num_of_channels, tensor_type - ); - } -#elif defined (HIP_COMPILE) - { - color_twist_hip_batch( - static_cast(srcPtr), - static_cast(dstPtr), - rpp::deref(rppHandle), - chn_format, num_of_channels - ); - } -#endif //BACKEND - - return RPP_SUCCESS; + return(color_twist_helper(RPPI_CHN_PACKED, 3, RPPTensorDataType::FP16, + srcPtr, srcSize, maxSrcSize, dstPtr, alpha, beta, + hueShift, saturationFactor, nbatchSize, rppHandle)); }