RPP Tensor Water Augmentation on HOST and HIP

include/rppt_tensor_effects_augmentations.h

@@ -272,6 +272,29 @@ RppStatus rppt_non_linear_blend_host(RppPtr_t srcPtr1, RppPtr_t srcPtr2, RpptDes
 RppStatus rppt_non_linear_blend_gpu(RppPtr_t srcPtr1, RppPtr_t srcPtr2, RpptDescPtr srcDescPtr, RppPtr_t dstPtr, RpptDescPtr dstDescPtr, Rpp32f *stdDevTensor, RpptROIPtr roiTensorPtrSrc, RpptRoiType roiType, rppHandle_t rppHandle);
 #endif // GPU_SUPPORT
 
+/******************** water ********************/
+
+// *param[in] srcPtr source tensor memory
+// *param[in] srcDescPtr source tensor descriptor
+// *param[out] dstPtr destination tensor memory
+// *param[in] dstDescPtr destination tensor descriptor
+// *param[in] amplitudeXTensor amplitudeX values for water effect (1D tensor of size batchSize)
+// *param[in] amplitudeYTensor amplitudeY values for water effect (1D tensor of size batchSize)
+// *param[in] freqXTensor freqX values for water effect (1D tensor of size batchSize)
+// *param[in] freqYTensor freqY values for water effect (1D tensor of size batchSize)
+// *param[in] phaseXTensor amplitudeY values for water effect (1D tensor of size batchSize)
+// *param[in] phaseYTensor amplitudeY values for water effect (1D tensor of size batchSize)
+// *param[in] roiTensorSrc ROI data for each image in source tensor (2D tensor of size batchSize * 4, in either format - XYWH(xy.x, xy.y, roiWidth, roiHeight) or LTRB(lt.x, lt.y, rb.x, rb.y))
+// *param[in] roiType ROI type used (RpptRoiType::XYWH or RpptRoiType::LTRB)
+// *returns a  RppStatus enumeration.
+// *retval RPP_SUCCESS : succesful completion
+// *retval RPP_ERROR : Error
+
+RppStatus rppt_water_host(RppPtr_t srcPtr, RpptDescPtr srcDescPtr, RppPtr_t dstPtr, RpptDescPtr dstDescPtr, Rpp32f *amplitudeXTensor, Rpp32f *amplitudeYTensor, Rpp32f *frequencyXTensor, Rpp32f *frequencyYTensor, Rpp32f *phaseXTensor, Rpp32f *phaseYTensor, RpptROIPtr roiTensorPtrSrc, RpptRoiType roiType, rppHandle_t rppHandle);
+#ifdef GPU_SUPPORT
+RppStatus rppt_water_gpu(RppPtr_t srcPtr, RpptDescPtr srcDescPtr, RppPtr_t dstPtr, RpptDescPtr dstDescPtr, Rpp32f *amplitudeXTensor, Rpp32f *amplitudeYTensor, Rpp32f *frequencyXTensor, Rpp32f *frequencyYTensor, Rpp32f *phaseXTensor, Rpp32f *phaseYTensor, RpptROIPtr roiTensorPtrSrc, RpptRoiType roiType, rppHandle_t rppHandle);
+#endif // GPU_SUPPORT
+
 #ifdef __cplusplus
 }
 #endif

src/include/cpu/rpp_cpu_common.hpp

@@ -5044,6 +5044,20 @@ inline void compute_generic_nn_srclocs_and_validate_sse(__m128 pSrcY, __m128 pSr
     _mm_storeu_si128((__m128i*) srcLoc, pxSrcLoc);
 }
 
+inline void compute_generic_nn_srclocs_and_validate_avx(__m256 pSrcY, __m256 pSrcX, __m256 *pRoiLTRB, __m256 pSrcStrideH, Rpp32s *srcLoc, Rpp32s *invalidLoad, bool hasRGBChannels = false)
+{
+    pSrcY = _mm256_round_ps(pSrcY, (_MM_FROUND_TO_NEAREST_INT |_MM_FROUND_NO_EXC));              // Nearest Neighbor Y location vector
+    pSrcX = _mm256_round_ps(pSrcX, (_MM_FROUND_TO_NEAREST_INT |_MM_FROUND_NO_EXC));              // Nearest Neighbor X location vector
+    _mm256_storeu_si256((__m256i*) invalidLoad, _mm256_cvtps_epi32(_mm256_or_ps(                 // Vectorized ROI boundary check
+        _mm256_or_ps(_mm256_cmp_ps(pSrcX, pRoiLTRB[0], _CMP_LT_OQ), _mm256_cmp_ps(pSrcY, pRoiLTRB[1],_CMP_LT_OQ)),
+        _mm256_or_ps(_mm256_cmp_ps(pSrcX, pRoiLTRB[2], _CMP_GT_OQ), _mm256_cmp_ps(pSrcY, pRoiLTRB[3], _CMP_GT_OQ))
+    )));
+    if (hasRGBChannels)
+        pSrcX = _mm256_mul_ps(pSrcX, avx_p3);
+    __m256i pxSrcLoc = _mm256_cvtps_epi32(_mm256_fmadd_ps(pSrcY, pSrcStrideH, pSrcX));
+    _mm256_storeu_si256((__m256i*) srcLoc, pxSrcLoc);
+}
+
 template <typename T>
 inline void compute_generic_nn_interpolation_pkd3_to_pln3(Rpp32f srcY, Rpp32f srcX, RpptROI *roiLTRB, T *dstPtrTempR, T *dstPtrTempG, T *dstPtrTempB, T *srcPtrChannel, RpptDescPtr srcDescPtr)
 {

src/modules/cpu/host_tensor_effects_augmentations.hpp

@@ -29,5 +29,6 @@ THE SOFTWARE.
 #include "kernel/noise_shot.hpp"
 #include "kernel/noise_gaussian.hpp"
 #include "kernel/non_linear_blend.hpp"
+#include "kernel/water.hpp"
 
 #endif // HOST_TENSOR_EFFECTS_AUGMENTATIONS_HPP

src/modules/hip/hip_tensor_effects_augmentations.hpp

@@ -29,5 +29,6 @@ THE SOFTWARE.
 #include "kernel/noise_shot.hpp"
 #include "kernel/noise_gaussian.hpp"
 #include "kernel/non_linear_blend.hpp"
+#include "kernel/water.hpp"
 
 #endif // HIP_TENSOR_EFFECTS_AUGMENTATIONS_HPP

src/modules/rppt_tensor_effects_augmentations.cpp

@@ -505,6 +505,95 @@ RppStatus rppt_non_linear_blend_host(RppPtr_t srcPtr1,
     return RPP_SUCCESS;
 }
 
+/******************** water ********************/
+
+RppStatus rppt_water_host(RppPtr_t srcPtr,
+                          RpptDescPtr srcDescPtr,
+                          RppPtr_t dstPtr,
+                          RpptDescPtr dstDescPtr,
+                          Rpp32f *amplitudeXTensor,
+                          Rpp32f *amplitudeYTensor,
+                          Rpp32f *frequencyXTensor,
+                          Rpp32f *frequencyYTensor,
+                          Rpp32f *phaseXTensor,
+                          Rpp32f *phaseYTensor,
+                          RpptROIPtr roiTensorPtrSrc,
+                          RpptRoiType roiType,
+                          rppHandle_t rppHandle)
+{
+    RppLayoutParams layoutParams = get_layout_params(srcDescPtr->layout, srcDescPtr->c);
+    if ((srcDescPtr->dataType == RpptDataType::U8) && (dstDescPtr->dataType == RpptDataType::U8))
+    {
+        water_u8_u8_host_tensor(static_cast<Rpp8u*>(srcPtr) + srcDescPtr->offsetInBytes,
+                                srcDescPtr,
+                                static_cast<Rpp8u*>(dstPtr) + dstDescPtr->offsetInBytes,
+                                dstDescPtr,
+                                amplitudeXTensor,
+                                amplitudeYTensor,
+                                frequencyXTensor,
+                                frequencyYTensor,
+                                phaseXTensor,
+                                phaseYTensor,
+                                roiTensorPtrSrc,
+                                roiType,
+                                layoutParams,
+                                rpp::deref(rppHandle));
+    }
+    else if ((srcDescPtr->dataType == RpptDataType::F16) && (dstDescPtr->dataType == RpptDataType::F16))
+    {
+        water_f16_f16_host_tensor(reinterpret_cast<Rpp16f*>(static_cast<Rpp8u*>(srcPtr) + srcDescPtr->offsetInBytes),
+                                  srcDescPtr,
+                                  reinterpret_cast<Rpp16f*>(static_cast<Rpp8u*>(dstPtr) + dstDescPtr->offsetInBytes),
+                                  dstDescPtr,
+                                  amplitudeXTensor,
+                                  amplitudeYTensor,
+                                  frequencyXTensor,
+                                  frequencyYTensor,
+                                  phaseXTensor,
+                                  phaseYTensor,
+                                  roiTensorPtrSrc,
+                                  roiType,
+                                  layoutParams,
+                                  rpp::deref(rppHandle));
+    }
+    else if ((srcDescPtr->dataType == RpptDataType::F32) && (dstDescPtr->dataType == RpptDataType::F32))
+    {
+        water_f32_f32_host_tensor(reinterpret_cast<Rpp32f*>(static_cast<Rpp8u*>(srcPtr) + srcDescPtr->offsetInBytes),
+                                  srcDescPtr,
+                                  reinterpret_cast<Rpp32f*>(static_cast<Rpp8u*>(dstPtr) + dstDescPtr->offsetInBytes),
+                                  dstDescPtr,
+                                  amplitudeXTensor,
+                                  amplitudeYTensor,
+                                  frequencyXTensor,
+                                  frequencyYTensor,
+                                  phaseXTensor,
+                                  phaseYTensor,
+                                  roiTensorPtrSrc,
+                                  roiType,
+                                  layoutParams,
+                                  rpp::deref(rppHandle));
+    }
+    else if ((srcDescPtr->dataType == RpptDataType::I8) && (dstDescPtr->dataType == RpptDataType::I8))
+    {
+        water_i8_i8_host_tensor(static_cast<Rpp8s*>(srcPtr) + srcDescPtr->offsetInBytes,
+                                srcDescPtr,
+                                static_cast<Rpp8s*>(dstPtr) + dstDescPtr->offsetInBytes,
+                                dstDescPtr,
+                                amplitudeXTensor,
+                                amplitudeYTensor,
+                                frequencyXTensor,
+                                frequencyYTensor,
+                                phaseXTensor,
+                                phaseYTensor,
+                                roiTensorPtrSrc,
+                                roiType,
+                                layoutParams,
+                                rpp::deref(rppHandle));
+    }
+
+    return RPP_SUCCESS;
+}
+
 /********************************************************************************************************************/
 /*********************************************** RPP_GPU_SUPPORT = ON ***********************************************/
 /********************************************************************************************************************/
@@ -1013,4 +1102,74 @@ RppStatus rppt_non_linear_blend_gpu(RppPtr_t srcPtr1,
 #endif // backend
 }
 
+RppStatus rppt_water_gpu(RppPtr_t srcPtr,
+                         RpptDescPtr srcDescPtr,
+                         RppPtr_t dstPtr,
+                         RpptDescPtr dstDescPtr,
+                         Rpp32f *amplitudeXTensor,
+                         Rpp32f *amplitudeYTensor,
+                         Rpp32f *frequencyXTensor,
+                         Rpp32f *frequencyYTensor,
+                         Rpp32f *phaseXTensor,
+                         Rpp32f *phaseYTensor,
+                         RpptROIPtr roiTensorPtrSrc,
+                         RpptRoiType roiType,
+                         rppHandle_t rppHandle)
+{
+#ifdef HIP_COMPILE
+    Rpp32u paramIndex = 0;
+    copy_param_float(amplitudeXTensor, rpp::deref(rppHandle), paramIndex++);
+    copy_param_float(amplitudeYTensor, rpp::deref(rppHandle), paramIndex++);
+    copy_param_float(frequencyXTensor, rpp::deref(rppHandle), paramIndex++);
+    copy_param_float(frequencyYTensor, rpp::deref(rppHandle), paramIndex++);
+    copy_param_float(phaseXTensor, rpp::deref(rppHandle), paramIndex++);
+    copy_param_float(phaseYTensor, rpp::deref(rppHandle), paramIndex);
+
+    if ((srcDescPtr->dataType == RpptDataType::U8) && (dstDescPtr->dataType == RpptDataType::U8))
+    {
+        hip_exec_water_tensor(static_cast<Rpp8u*>(srcPtr) + srcDescPtr->offsetInBytes,
+                              srcDescPtr,
+                              static_cast<Rpp8u*>(dstPtr) + dstDescPtr->offsetInBytes,
+                              dstDescPtr,
+                              roiTensorPtrSrc,
+                              roiType,
+                              rpp::deref(rppHandle));
+    }
+    else if ((srcDescPtr->dataType == RpptDataType::F16) && (dstDescPtr->dataType == RpptDataType::F16))
+    {
+        hip_exec_water_tensor(reinterpret_cast<half*>(static_cast<Rpp8u*>(srcPtr) + srcDescPtr->offsetInBytes),
+                              srcDescPtr,
+                              reinterpret_cast<half*>(static_cast<Rpp8u*>(dstPtr) + dstDescPtr->offsetInBytes),
+                              dstDescPtr,
+                              roiTensorPtrSrc,
+                              roiType,
+                              rpp::deref(rppHandle));
+    }
+    else if ((srcDescPtr->dataType == RpptDataType::F32) && (dstDescPtr->dataType == RpptDataType::F32))
+    {
+        hip_exec_water_tensor(reinterpret_cast<Rpp32f*>(static_cast<Rpp8u*>(srcPtr) + srcDescPtr->offsetInBytes),
+                              srcDescPtr,
+                              reinterpret_cast<Rpp32f*>(static_cast<Rpp8u*>(dstPtr) + dstDescPtr->offsetInBytes),
+                              dstDescPtr,
+                              roiTensorPtrSrc,
+                              roiType,
+                              rpp::deref(rppHandle));
+    }
+    else if ((srcDescPtr->dataType == RpptDataType::I8) && (dstDescPtr->dataType == RpptDataType::I8))
+    {
+        hip_exec_water_tensor(static_cast<Rpp8s*>(srcPtr) + srcDescPtr->offsetInBytes,
+                              srcDescPtr,
+                              static_cast<Rpp8s*>(dstPtr) + dstDescPtr->offsetInBytes,
+                              dstDescPtr,
+                              roiTensorPtrSrc,
+                              roiType,
+                              rpp::deref(rppHandle));
+    }
+
+    return RPP_SUCCESS;
+#elif defined(OCL_COMPILE)
+    return RPP_ERROR_NOT_IMPLEMENTED;
+#endif // backend
+}
+
 #endif // GPU_SUPPORT

utilities/test_suite/HIP/Tensor_hip.cpp

@@ -473,6 +473,35 @@ int main(int argc, char **argv)
 
                 break;
             }
+            case 29:
+            {
+                testCaseName = "water";
+
+                Rpp32f amplX[batchSize];
+                Rpp32f amplY[batchSize];
+                Rpp32f freqX[batchSize];
+                Rpp32f freqY[batchSize];
+                Rpp32f phaseX[batchSize];
+                Rpp32f phaseY[batchSize];
+
+                for (i = 0; i < batchSize; i++)
+                {
+                    amplX[i] = 2.0f;
+                    amplY[i] = 5.0f;
+                    freqX[i] = 5.8f;
+                    freqY[i] = 1.2f;
+                    phaseX[i] = 10.0f;
+                    phaseY[i] = 15.0f;
+                }
+
+                startWallTime = omp_get_wtime();
+                if (inputBitDepth == 0 || inputBitDepth == 1 || inputBitDepth == 2 || inputBitDepth == 5)
+                    rppt_water_gpu(d_input, srcDescPtr, d_output, dstDescPtr, amplX, amplY, freqX, freqY, phaseX, phaseY, roiTensorPtrSrc, roiTypeSrc, handle);
+                else
+                    missingFuncFlag = 1;
+
+                break;
+            }
             case 31:
             {
                 testCaseName = "color_cast";

utilities/test_suite/HIP/runTests.py

@@ -388,14 +388,17 @@ def rpp_test_suite_parser_and_validator():
             print("Unable to open results in " + RESULTS_DIR + "/consolidated_results_" + TYPE + ".stats.csv")
 
 # print the results of qa tests
-supportedCaseList = ['0', '1', '2', '4', '13', '31', '34', '36', '37', '38', '54', '84', '87']
+supportedCaseList = ['0', '1', '2', '4', '13', '29', '31', '34', '36', '37', '38', '54', '84', '87']
 nonQACaseList = ['54', '84']
 supportedCases = 0
 for num in caseList:
     if qaMode == 1 and num not in nonQACaseList:
         supportedCases += 1
     elif qaMode == 0 and num in supportedCaseList:
         supportedCases += 1
+    elif qaMode == 0 and num in supportedCaseList:
+        supportedCases += 1
+
 caseInfo = "Tests are run for " + str(supportedCases) + " supported cases out of the " + str(len(caseList)) + " cases requested"
 if qaMode and testType == 0:
     qaFilePath = os.path.join(outFilePath, "QA_results.txt")

utilities/test_suite/HOST/Tensor_host.cpp

@@ -454,6 +454,35 @@ int main(int argc, char **argv)
 
                     break;
                 }
+                case 29:
+                {
+                    testCaseName = "water";
+
+                    Rpp32f amplX[batchSize];
+                    Rpp32f amplY[batchSize];
+                    Rpp32f freqX[batchSize];
+                    Rpp32f freqY[batchSize];
+                    Rpp32f phaseX[batchSize];
+                    Rpp32f phaseY[batchSize];
+
+                    for (i = 0; i < batchSize; i++)
+                    {
+                        amplX[i] = 2.0f;
+                        amplY[i] = 5.0f;
+                        freqX[i] = 5.8f;
+                        freqY[i] = 1.2f;
+                        phaseX[i] = 10.0f;
+                        phaseY[i] = 15.0f;
+                    }
+
+                    startWallTime = omp_get_wtime();
+                    if (inputBitDepth == 0 || inputBitDepth == 1 || inputBitDepth == 2 || inputBitDepth == 5)
+                        rppt_water_host(input, srcDescPtr, output, dstDescPtr, amplX, amplY, freqX, freqY, phaseX, phaseY, roiTensorPtrSrc, roiTypeSrc, handle);
+                    else
+                        missingFuncFlag = 1;
+
+                    break;
+                }
                 case 31:
                 {
                     testCaseName = "color_cast";

utilities/test_suite/HOST/runTests.py

@@ -162,14 +162,17 @@ def rpp_test_suite_parser_and_validator():
 subprocess.call(["./testAllScript.sh", srcPath1, args.input_path2, str(testType), str(numRuns), str(qaMode), str(decoderType), str(preserveOutput), str(batchSize), " ".join(caseList)])  # nosec
 
 # print the results of qa tests
-supportedCaseList = ['0', '1', '2', '4', '13', '31', '34', '36', '37', '38', '84', '87']
+supportedCaseList = ['0', '1', '2', '4', '13', '29', '31', '34', '36', '37', '38', '84', '87']
 nonQACaseList = ['54', '84']
 supportedCases = 0
 for num in caseList:
     if qaMode == 1 and num not in nonQACaseList:
         supportedCases += 1
     elif qaMode == 0 and num in supportedCaseList:
         supportedCases += 1
+    elif qaMode == 0 and num in supportedCaseList:
+        supportedCases += 1
+
 caseInfo = "Tests are run for " + str(supportedCases) + " supported cases out of the " + str(len(caseList)) + " cases requested"
 if qaMode and testType == 0:
     qaFilePath = os.path.join(outFilePath, "QA_results.txt")

utilities/test_suite/rpp_test_suite_common.h

@@ -64,6 +64,7 @@ std::map<int, string> augmentationMap =
     {2, "blend"},
     {4, "contrast"},
     {13, "exposure"},
+    {29, "water"},
     {31, "color_cast"},
     {34, "lut"},
     {36, "color_twist"},