Audio HIP PR1 - Non Silent Region HIP Support

[sampath1117]

• Adds support for Non Silent Region in HIP Backend
• Adds relevant qa tests and performance test support

[snehaa8]

@sampath1117 Please address minor comments

[snehaa8]

Why are we initially setting shared memory as char then reinterpreting to float?

[sampath1117]

Removed it. Used a direct declaration

[snehaa8]

Please rename this and use it in kernel executor as ...Arr

[snehaa8]

Please remove extra line here

[snehaa8]

Its always good to have new line at end of file

[sampath1117]

@r-abishek
I have addressed all review comments

[r-abishek]

Pls address comments

[r-abishek]

You can directly do a return hip_exec_*();

[r-abishek]

Lets remove the print statements from lib code on L334 and L339.
Between the error code here, its definition in rppdefs, and any restrictions for this function mentioned in the API header file, the user should be able to figure out the issue.
Also, put the remaining code in an 'else', where the else part alone returns SUCCESS.

[r-abishek]

The protocol on this naming discussed was to use camelCase_smem only if the variable is shared memory. Naming functions and local variables as smem too is a little confusing. Maybe call the function "compute_position_in_smem()"

Also suggest vectorizing variables to pos_i2.x and pos_i2.y, and use an offset_i2 instead of offsetA and offsetB

[r-abishek]

You can directly do int twiceThreadID = 2 * hipThreadIdx_x; since tid is not really used, you seem to only use 2 * tid

[r-abishek]

And on L32, you can use if (hipThreadIdx_x == 0) directly

[r-abishek]

Can possibly do something like this:

    for (int d = bufferLength >> 1; d > 0; d >>= 1)
    {
        __syncthreads();
        int dMul2 = 2 * d * offset;
        int increment = blockDimMul2 * offset;
        for (int idxMul2 = twiceThreadID * offset; idxMul2 < dMul2; idxMul2 += increment)
        {
            int pos_i2 = (make_int2(idxMul2) + offset_i2) / 32;     // set a constant int2 = 32 in advance
            input[pos_i2.y] += input[pos_i2.x];
        }
        offset <<= 1;
        offset_i2 = make_int2(offset - 1, 2 * offset - 1);
    }

[sampath1117]

@r-abishek
I have addressed all comments and also removed the hipMalloc and hipFree for MMS buffer inside kernel and used handle memory instead and it was working fine

[r-abishek]

@sampath1117 Added some minor comments

[r-abishek]

If error code 20, 21, 22 are like 14, pls use the word "insufficient" like 14 instead of saying "out of bound"

[sampath1117]

Checked 14th error code, RPP_ERROR_INSUFFICIENT_DST_BUFFER_LENGTH
It is not like 14

In NSR HIP kernel case, since we are using pre-allocated scratch memory for some intermediate buffers and for a certain input configuration if the required scratch memory goes beyond the bounds, we return this error codes

Where as insufficient are the cases, where the user passed buffers are lesser than required memory
https://github.com/sampath1117/rpp/blob/sr/nsr_hip/src/modules/rppt_tensor_statistical_operations.cpp#L123

So not changing the names here for new error codes added

[r-abishek]

Suggest stating what these numbers are - 600000, 293, 128

[sampath1117]

600000 is the maximum size that will be required for MMS buffer based on RNNT dataset
293 is the size required for storing reduction outputs for 600000 sample
128 is the size required for storing cutoff values for batch size 128

Added this details

[r-abishek]

This device helper function could use some comments

[sampath1117]

Done
Somehow it is not getting reflected in github UI
https://github.com/sampath1117/rpp/blob/sr/nsr_hip/src/modules/hip/kernel/non_silent_region_detection.hpp#L14

[r-abishek]

@sampath1117 Please pull latest develop into this branch again, ideally the changes from this PR for the audio support flag should show - https://github.com/ROCm/rpp/pull/378/files#diff-60d74aa0b8e5fb394c2c48769bcd9be4c9481762d535f37eca99b957e0440f3fR122-R125

[r-abishek]

Pls check other changes too in the test_suite, and test_suite/HIP CMake - https://github.com/ROCm/rpp/pull/378/files#diff-60d74aa0b8e5fb394c2c48769bcd9be4c9481762d535f37eca99b957e0440f3fR70-R77

[sampath1117]

Added the required changes in HIP test suite CMake