RPP Rain Tensor

[HazarathKumarM]

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

[sampath1117]

@HazarathKumarM
Please address first round of review comments

[sampath1117]

I think most of the inputs, outputs are stored in .png format in the doxygen folder
Here the input should be img150x150.png instead of img150x150.jpg

Please save effects_augmentations_rain_img150x150.jpg also in .png format for consistency
Ensure to make this change for both HOST and HIP

[HazarathKumarM]

Done

[sampath1117]

It might better to give some ranges for rainHeight in documentation
since we giving rainHeight as height/2 also does give any meaninful output

Please add some documentation like 10% of input height or 1% of input height like that
You need to do the experiments and then only arrive at this number

[sampath1117]

Same comment
Give some range of values for slant which gives meaningful output

[sampath1117]

Also for now
Please clear mention in documentation like
Rain function will work only when input images are of same size

[HazarathKumarM]

Done

[sampath1117]

You can probably assign initValue as 0 during declaration itself and use only an if condition for Rpp8s case

[HazarathKumarM]

Done

[sampath1117]

Same comment as in HOST kernel
slant / rainHeight is a const value, need not recompute

[HazarathKumarM]

DOne

[sampath1117]

Please refer to comments added in HOST kernel for this function

[HazarathKumarM]

Done

[sampath1117]

For batchsize
Please ensure to get the value from dstDescPtr->n instead of handle.GetBatchSize() function as discussed before

[HazarathKumarM]

Done

[sampath1117]

I think it is not advisable to use pinned memory for mask here
since this can be a big buffer based on input

Would suggest to compute this mask in HOST scratch buffer and copy this to HIP scratch buffer
Try an experiment with this and see if there is any performance improvement (with batchsize 128)

[HazarathKumarM]

Done

[sampath1117]

@HazarathKumarM
Please address review comments

[sampath1117]

we can move compute_rain_32_host() before compute_rain_24_host()

[HazarathKumarM]

Done

[sampath1117]

Accept pMul as a reference variable. Applicable to all such functions added

Suggested change

	inline void compute_rain_48_host(__m256 *p1, __m256 *p2, __m256 pMul)
	inline void compute_rain_48_host(__m256 *p1, __m256 *p2, __m256 &pMul)

[HazarathKumarM]

Done

[sampath1117]

@HazarathKumarM
is this addressed, chrono removal?

[sampath1117]

Please change int k to Rpp32u k

Suggested change

	for (int k = 0; k < rainWidth; k++)
	for (Rpp32u k = 0; k < rainWidth; k++)

[HazarathKumarM]

Done

[sampath1117]

Do we still need k * srcDescPtr->strides.wStride multiplication here?
since it is just a grey scale mask, it should be ideally like below?

Suggested change

	*(rainLayerTemp + k * srcDescPtr->strides.wStride) = rainValue;
	rainLayerTemp[k] = rainValue;

[HazarathKumarM]

Done

[sampath1117]

Please do memset only for w * h values as mentioned in CPU kernel comments

[HazarathKumarM]

Done

[sampath1117]

Lets change globalThreads_x as below

Suggested change

	int globalThreads_x = (dstDescPtr->strides.hStride + 7) >> 3;
	int globalThreads_x = (dstDescPtr->w + 7) >> 3;

[HazarathKumarM]

Done

[sampath1117]

Remove this modification at L242

[HazarathKumarM]

Done

[sampath1117]

Please rename the name as srcPtr instead of srcPtr1

Suggested change

	RppStatus hip_exec_rain_tensor(T *srcPtr1,
	RppStatus hip_exec_rain_tensor(T *srcPtr

[HazarathKumarM]

Done

[sampath1117]

Lets remove this modification, since this is not needed

[HazarathKumarM]

Done

[r-abishek]

@HazarathKumarM Please address comments

[r-abishek]

Innstead of the 150x150, please use one of the larger dashcam images here for all 3 - fog, rain and snow, and ensure the most realistic output is used as the documentation image

[HazarathKumarM]

Done

[r-abishek]

Please add "To be tuned by user depending on size of image" for the rainWidth and rainHeight

[HazarathKumarM]

Done

[r-abishek]

Can the user give something like -400 for slant? Please mention unit, range, datatype. I see the unit is pixels, but based on the one line of information provided in the docs, its hard to understand how slant can be defined in pixels. People might expect slant in degrees. Say a 30deg or a -15deg slant.

Something like "Slant of the rain drops (A single Rpp32s/f representing slant of raindrops in degrees - Values in range [-90, 90] representing extremes of right slant to left slant)"

Also I notice the docs say negative or positive values, but the API only takes Rpp32u (which is unsigned int) and that's a bug. Please do re-test with negative inputs too

[HazarathKumarM]

Done

[r-abishek]

What's the range of alpha values? We can define it either based on realism of image or max/min values possible. Can it be 0 or negative?

Please refer to say brightness or blend for alpha blending for the level of explanation needed - both have different ranges:
brightness -> https://github.com/ROCm/rpp/blob/develop/include/rppt_tensor_color_augmentations.h#L55
blend -> https://github.com/ROCm/rpp/blob/develop/include/rppt_tensor_color_augmentations.h#L143

[HazarathKumarM]

DOne

[r-abishek]

State why in an inline comment. What does the magic number 0.004 do?

[HazarathKumarM]

Done

[r-abishek]

I really think giving an angle instead might be more convenient/intuitive to the end user. Pls check.

[HazarathKumarM]

Done

[r-abishek]

Our scratchBufferPinned seems to be decently large. It can hold a 1 channel 4K float image. Please substitute L202 to instead point to scratchBufferPinned.
And then you can remove L230, L231 and directly pass the pinned memory buffer pointer, rainLayer into the kernel.
Pls recheck outputs.

[HazarathKumarM]

If we use scratchPinned Buffer we observed very high times and performance is very low since the buffer size huge.

with the current config I got more than 2x improvement compared to the pinned memory

[r-abishek]

Are we sure we are not re-adding these loads in this file? Always check if they already exist

[HazarathKumarM]

verified

[r-abishek]

Why the magic number? Pls add a comment

[HazarathKumarM]

Done

[r-abishek]

What percentage of time do just these 3 lines L86-L88 take as compared to the whole function?

[HazarathKumarM]

around 50-60% on average