RPP Tensor Support - Snow on HOST and HIP #347

Dineshbabu-Ravichandran · 2024-09-25T04:26:13Z

Adds tensor support for Snow Augmentation optimized using AVX2 on HOST backend
Adds tensor support for Snow Augmentation on HIP backend
Adds unit and performance tests support for the Snow Augmentation in test suite

Srihari-mcw · 2024-09-26T01:10:10Z

I think this output should be on the unit tests default 150x150 image. Pls check once @Dineshbabu-Ravichandran

Srihari-mcw · 2024-09-26T01:13:01Z

include/rppt_tensor_effects_augmentations.h

+ * \details The Snow augmentation does a  modification of brightness on a batch of RGB(3 channel) / greyscale(1 channel) images with an NHWC/NCHW tensor layout.<br>
+ * - srcPtr depth ranges - Rpp8u (0 to 255), Rpp16f (0 to 1), Rpp32f (0 to 1), Rpp8s (-128 to 127).
+ * - dstPtr depth ranges - Will be same depth as srcPtr.
+ * \image html img640x480.png Sample Input


Same doubt here

We used this dashcam image for fog PR . https://github.com/r-abishek/rpp/pull/332/files#diff-006ea80a28a2d71eeb553a7a9c8b32912f4d35f871150a5094fbe85e3503575f . So I used same .

Srihari-mcw · 2024-09-26T01:15:20Z

include/rppt_tensor_effects_augmentations.h

+ * \param [in] brightnessCoefficient brightness modification parameter for snow calculation (1D tensor in HOST memory, of size batchSize with 1 < brightnessCoefficient[i] <= 4 for each image in batch)
+ * \param [in] snowThreshold threshold parameter for snow calculation (1D tensor in HOST memory, of size batchSize with 0 < snowThresholdTensor[i] <= 1 for each image in batch)
+ * \param [in] darkMode darkMode  values to set dark mode on/off (1D tensor in HOST memory, of size batchSize, with darkModeTensor[i] = 0/1)
+ * \param [in] roiTensorSrc ROI data in HOST memory, 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))


Change to roiTensorPtrSrc

Srihari-mcw · 2024-09-26T01:16:53Z

include/rppt_tensor_effects_augmentations.h

+ * \param [in] brightnessCoefficient brightness modification parameter for snow calculation (1D tensor in pinned/HIP memory, of size batchSize with 1 < brightnessCoefficient[i] <= 4 for each image in batch)
+ * \param [in] snowThreshold threshold parameter for snow calculation (1D tensor in pinned/HIP memory, of size batchSize with 0 < snowThreshold[i] <= 1 for each image in batch)
+ * \param [in] darkMode darkMode  values to set dark mode on/off (1D tensor in pinned/HIP memory, of size batchSize, with darkModeTensor[i] = 0/1)
+ * \param [in] roiTensorSrc ROI data in HIP memory, 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))


Change to roiTensorPtrSrc

Srihari-mcw · 2024-09-26T01:21:15Z

src/include/cpu/rpp_cpu_common.hpp

@@ -3135,6 +3138,217 @@ inline void compute_color_temperature_24_host(__m256 *p, __m256 pAdj)
    p[2] = _mm256_sub_ps(p[2], pAdj);    // color_temperature adjustment Bs
 }

+inline void compute_snow_host(RpptFloatRGB *pixel, Rpp32f brightnessCoefficient, Rpp32f snowCoefficient, Rpp32s darkMode)


These functions cannot be part of the host code for snow itself snow.hpp @Dineshbabu-Ravichandran @sampath1117 Pls share your thoughts

Because in warp perspective case I remember the helper functions to be part of same host file

Srihari-mcw · 2024-09-26T01:25:48Z

src/include/cpu/rpp_cpu_common.hpp

+    pH = avx_p0;                                                                                                            // hue = 0.0f;
+    pS = avx_p0;                                                                                                            // sat = 0.0f;
+    pAdd = avx_p0;                                                                                                          // add = 0.0f;
+    pL = _mm256_mul_ps(_mm256_add_ps(pCmax, pCmin), _mm256_set1_ps(0.5f));                                                   //  l = delta * 0.5


I think the comment is wrong in this line

Srihari-mcw · 2024-09-26T01:36:39Z

src/include/cpu/rpp_cpu_simd.hpp

@@ -1555,6 +1574,13 @@ inline void rpp_load24_f32pln3_to_f32pln3_avx(Rpp32f *srcPtrR, Rpp32f *srcPtrG,
    p[2] = _mm256_loadu_ps(srcPtrB);
 }

+inline void rpp_load24_f16pln3_to_f32pln3_avx(Rpp16f *srcPtrR, Rpp16f *srcPtrG, Rpp16f *srcPtrB, __m256 *p)


Pls move this function before rpp_load24_f16pln3_to_f32pln3_avx as in threshold implementation

https://github.com/r-abishek/rpp/pull/322/files#diff-ab58bec6af3335be388abefa97eab0a886da47cdb1afb124e60a1a94de14c7b5

Srihari-mcw · 2024-09-26T01:37:28Z

src/include/cpu/rpp_cpu_simd.hpp

@@ -1647,6 +1673,11 @@ inline void rpp_load8_f32_to_f32_avx(Rpp32f *srcPtr, __m256 *p)
    p[0] = _mm256_loadu_ps(srcPtr);
 }

+inline void rpp_load8_f16_to_f32_avx(Rpp16f *srcPtr, __m256 *p)


Pls move this function before rpp_load8_f32_to_f64_avx

Srihari-mcw · 2024-09-26T02:10:13Z

src/modules/hip/kernel/snow.hpp

+
+    int globalThreads_x = (dstDescPtr->strides.hStride + 7) >> 3;
+    int globalThreads_y = dstDescPtr->h;
+    int globalThreads_z = handle.GetBatchSize();


int globalThreads_z = dstDescPtr->n; maybe could be used here?

Srihari-mcw · 2024-09-26T02:14:26Z

src/modules/rppt_tensor_effects_augmentations.cpp

+    }
+    else if ((srcDescPtr->dataType == RpptDataType::F16) && (dstDescPtr->dataType == RpptDataType::F16))
+    {
+        snow_f16_f16_host_tensor((Rpp16f*) (static_cast<Rpp8u*>(srcPtr) + srcDescPtr->offsetInBytes),


Pls do reinterpret_cast<Rpp16f*> here

Srihari-mcw · 2024-09-26T02:14:46Z

src/modules/rppt_tensor_effects_augmentations.cpp

+    }
+    else if ((srcDescPtr->dataType == RpptDataType::F32) && (dstDescPtr->dataType == RpptDataType::F32))
+    {
+        snow_f32_f32_host_tensor((Rpp32f*) (static_cast<Rpp8u*>(srcPtr) + srcDescPtr->offsetInBytes),


Pls do reinterpret_cast<Rpp32f*> here

Srihari-mcw · 2024-09-26T02:15:41Z

src/modules/rppt_tensor_effects_augmentations.cpp

+    }
+    else if ((srcDescPtr->dataType == RpptDataType::F16) && (dstDescPtr->dataType == RpptDataType::F16))
+    {
+        hip_exec_snow_tensor((half*) (static_cast<Rpp8u*>(srcPtr) + srcDescPtr->offsetInBytes),


reinterpret_cast<half*> here

Srihari-mcw · 2024-09-26T02:16:00Z

src/modules/rppt_tensor_effects_augmentations.cpp

+    }
+    else if ((srcDescPtr->dataType == RpptDataType::F32) && (dstDescPtr->dataType == RpptDataType::F32))
+    {
+        hip_exec_snow_tensor((Rpp32f*) (static_cast<Rpp8u*>(srcPtr) + srcDescPtr->offsetInBytes),


Use reinerpret_cast

HazarathKumarM

@Dineshbabu-Ravichandran please resolve the comments

HazarathKumarM · 2024-09-26T08:52:18Z

src/modules/hip/kernel/snow.hpp

+        }
+        else if ((srcDescPtr->layout == RpptLayout::NCHW) && (dstDescPtr->layout == RpptLayout::NHWC))
+        {
+            globalThreads_x = (srcDescPtr->strides.hStride + 7) >> 3;


I think this line is repeated , the same code is there in L347

I removed the line L406.

HazarathKumarM · 2024-09-26T08:55:36Z

src/modules/hip/kernel/snow.hpp

+    if (roiType == RpptRoiType::LTRB)
+        hip_exec_roi_converison_ltrb_to_xywh(roiTensorPtrSrc, handle);
+
+


remove the empty line here

HazarathKumarM · 2024-09-26T08:56:00Z

src/modules/hip/kernel/snow.hpp

+        rpp_hip_load8_and_unpack_to_float8(srcPtr + srcIdx, &pix_f8);
+        snow_hip_compute(srcPtr, &pix_f8, brightnessCoefficient, snowThreshold, darkMode);
+        rpp_hip_pack_float8_and_store8(dstPtr + dstIdx, &pix_f8);
+


remove the empty line

HazarathKumarM · 2024-09-26T08:57:56Z

src/modules/hip/kernel/snow.hpp

+
+__device__ __forceinline__ void snow_8RGB_hip_compute(d_float24 *pix_f24, float *brightnessCoefficient, float *snowThreshold, int *darkMode)
+{
+    snow_1RGB_hip_compute(&(pix_f24->f1[ 0]), &(pix_f24->f1[ 8]), &(pix_f24->f1[16]), brightnessCoefficient, snowThreshold, darkMode);


remove the extra space inside [] brackets

HazarathKumarM · 2024-09-26T09:00:01Z

src/modules/hip/kernel/snow.hpp

+
+__device__ __forceinline__ void snow_1GRAY_hip_compute(float *pixel, float *brightnessCoefficient, float *snowThreshold, int *darkMode)
+{
+    float l = *pixel;


please use some meaningful variable name instead of 'l'

I chnaged l to lightness.

HazarathKumarM · 2024-09-26T09:02:22Z

src/modules/hip/kernel/snow.hpp

+        l = l * fmaf((brightnessFactor - 1.0f), (1.0f - (l - lower_threshold) / (upper_threshold - lower_threshold)), 1.0f);
+    }
+    // Modify L 
+    if(l <= *snowThreshold)


remove the {} brackets for this IF statement

HazarathKumarM · 2024-09-26T09:05:07Z

src/include/cpu/rpp_cpu_common.hpp

+    pixel->R = hueCoefficient[0];
+    pixel->G = hueCoefficient[1];
+    pixel->B = hueCoefficient[2];
+


please the empty line here

HazarathKumarM · 2024-09-26T09:12:40Z

src/modules/hip/kernel/snow.hpp

+__device__ __forceinline__ void snow_1GRAY_hip_compute(float *pixel, float *brightnessCoefficient, float *snowThreshold, int *darkMode)
+{
+    float l = *pixel;
+    float lower_threshold = 0.0f;


Don't use snake case for variable names, use only camel case

HazarathKumarM · 2024-09-26T09:14:59Z

src/modules/hip/kernel/snow.hpp

+
+    // Modify L 
+    if(l <= *snowThreshold && !((hue >= 0.514f && hue <= 0.63f) && (sat >= 0.196f) && (l >= 0.196f)))
+    {


please remove {} brackets if there is only one line inside conditional statements or loops,

Please check all such instances in this PR and remove the brackets

HazarathKumarM · 2024-09-26T11:04:13Z

src/modules/hip/kernel/snow.hpp

+    *pixelG = rgb_f4.y;
+    *pixelB = rgb_f4.z;
+}
+


please remove the empty line here

* Make initial changes for raw CPP version of warp perspective * Fix calls to compute_warp_perspective_src_loc function * Update changes to go through nearest neighbours case * AVX HOST codes for warp perspective initial * Fixes for accuracy in warp perspective * More fixes for accuracy in warp perspective * Update the cide for AVX version of Planar to Planar * Add bilinear u8 host code for warp perspective * Make updates to include functions for F32 data type * Make updates to use cast instead of set and fix issues with raw C implementation * Add i8 host codes * Add updates for F16 Bilinear Code * Update the initial HIP code for warp perspective * Update fixes for HIP code * Add Warp Perspective Nearest Neighbors F16 code for PKD3_to_PLN3 and PLN3_to_PLN3 * Add updates for PLN to PLN configuration * Add updates for PKD3 to PKD3 case * Rename variables * Update changes to log images separately for Bilinear and Nearest Neighbors * fixed bug in raw c code of PKD-PLN variant * minor bug fix for F16 PLN variants * minor fixes in HOST test suite * Update the HIP code for review comments and refactoring of device functions * Update the comments alignment * Rename functions and add cases in HOST and HIP runTests.py * Update indentations for compuatations and rename vectors * Update documentations and add more reference variables * Make more formatting changes * Make further updates by including test cases * Make updates to use reinterpret cast * Update reinterpret casts for PLN to PLN configuration u8 and i8 codes * Make updates to enclose code inside AVX2 flag * Make further changes to update type casting * Update the version * Make updates to add warp perspective image * Modify comments, update CHANGELOG and update flags * Update further comments in warp perspective * Add more comments for warp perspective * Update based on further review comments * Update the case number for warp_perspective in common.py * Address review comments * Make initial changes for raw CPP version of warp perspective * Fix calls to compute_warp_perspective_src_loc function * Update changes to go through nearest neighbours case * AVX HOST codes for warp perspective initial * Fixes for accuracy in warp perspective * More fixes for accuracy in warp perspective * Update the cide for AVX version of Planar to Planar * Add bilinear u8 host code for warp perspective * Make updates to include functions for F32 data type * Make updates to use cast instead of set and fix issues with raw C implementation * Add i8 host codes * Add updates for F16 Bilinear Code * Update the initial HIP code for warp perspective * Update fixes for HIP code * Add Warp Perspective Nearest Neighbors F16 code for PKD3_to_PLN3 and PLN3_to_PLN3 * Add updates for PLN to PLN configuration * Add updates for PKD3 to PKD3 case * Rename variables * Update changes to log images separately for Bilinear and Nearest Neighbors * fixed bug in raw c code of PKD-PLN variant * minor bug fix for F16 PLN variants * minor fixes in HOST test suite * Update the HIP code for review comments and refactoring of device functions * Update the comments alignment * Rename functions and add cases in HOST and HIP runTests.py * Update indentations for compuatations and rename vectors * Update documentations and add more reference variables * Make more formatting changes * Make further updates by including test cases * Make updates to use reinterpret cast * Update reinterpret casts for PLN to PLN configuration u8 and i8 codes * Make updates to enclose code inside AVX2 flag * Make further changes to update type casting * Make updates to add warp perspective image * Modify comments, update CHANGELOG and update flags * Update further comments in warp perspective * Add more comments for warp perspective * Update based on further review comments * Update the case number for warp_perspective in common.py * Address review comments * Fix conflits with warp perspective * Update version details * Merge branch 'ar/opt_warp_perspective' of https://github.com/r-abishek/rpp into opt_warp_perspective_rebased * Update version to 1.9.10 including warp perspective * Updates to convert to XYWH from LTRB instead of opposite * Update CHANGELOG.md Co-authored-by: spolifroni-amd <Sandra.Polifroni@amd.com> * Revert changes and convert to ltrb instead of xywh --------- Co-authored-by: Srihari-mcw <srihari@multicorewareinc.com> Co-authored-by: sampath1117 <sampath.rachumallu@multicorewareinc.com> Co-authored-by: Kiriti Gowda <kiritigowda@gmail.com> Co-authored-by: Rajy Rawther <Rajy.MeeyakhanRawther@amd.com> Co-authored-by: spolifroni-amd <Sandra.Polifroni@amd.com>

Dineshbabu-Ravichandran and others added 24 commits September 9, 2024 14:25

Added Initial code for Snow HIP

6f48144

Merging with develop branch and resolve conflits

2170c24

Added brightness Coefficient Parameter

e0b1968

Changes in function definition

1ff68c9

Resolved the issue of black and white images

2983e56

Modified to handle float variant

72f53df

Handling dark images and Performance improvement

5e1ae7d

Optimized the conditional statement for HUE calculation

49c5f41

Reverting the changes in Hue calaculation due to performance issue

fa8603c

Performance improvement by removing fuction call for Hue to RGB

427658e

Modified the code to run on Single Thread

1144414

Reverted the approach of using single thread due to performance issue

30e3162

Implemented some initial Implementation on Host

e61cf29

Removing Print statements

5b1f462

Support for pln1 variant is added

cd44807

Merge branch 'develop' into db/opt_snow

c8d5644

Added support for F32,F16,I8 variant

98bd8b7

Added dark mode toggle for HIP side

f23fde2

Implementation of dark mode on HOST

1f5c163

Code cleanup

c019348

Changes in Description

224d516

Merge branch 'develop' into db/opt_snow

65a1f70

Code cleaned up and added Golden outputs

2f01a62

Performed some code optimization and reviewed the code

aaec5cd

Srihari-mcw reviewed Sep 26, 2024

View reviewed change 8000 s

Srihari-mcw reviewed Sep 26, 2024

View reviewed changes

Changes done based on review Comments

bfa10d6

Dineshbabu-Ravichandran changed the base branch from develop to master September 26, 2024 11:13

Dineshbabu-Ravichandran changed the base branch from master to develop September 26, 2024 11:13

HazarathKumarM reviewed Sep 26, 2024

View reviewed changes

Dineshbabu-Ravichandran and others added 8 commits September 26, 2024 11:46

AVX code block is changed

e0b1851

Chnages made based on Review Commands

99cc08d

Minor changes on variable name

a294fc9

Doxygen Image changed

aa300b0

Merge branch 'develop' into db/opt_snow

ed2346e

Merge branch 'develop' into opt_snow

503f71c

Minor fix due to merging

e2e23b9

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

RPP Tensor Support - Snow on HOST and HIP #347

RPP Tensor Support - Snow on HOST and HIP #347

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

		if (roiType == RpptRoiType::LTRB)
		hip_exec_roi_converison_ltrb_to_xywh(roiTensorPtrSrc, handle);

RPP Tensor Support - Snow on HOST and HIP #347

Are you sure you want to change the base?

RPP Tensor Support - Snow on HOST and HIP #347

Uh oh!

Conversation

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!