Audio HIP PR3 - Downmixing HIP Support #261
There are no files selected for viewing
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,30 @@ | ||
| /* | ||
| MIT License | ||
|
|
||
| Copyright (c) 2019 - 2024 Advanced Micro Devices, Inc. | ||
|
|
||
| Permission is hereby granted, free of charge, to any person obtaining a copy | ||
| of this software and associated documentation files (the "Software"), to deal | ||
| in the Software without restriction, including without limitation the rights | ||
| to use, copy, modify, merge, publish, distribute, sublicense, and/or sell | ||
| copies of the Software, and to permit persons to whom the Software is | ||
| furnished to do so, subject to the following conditions: | ||
|
|
||
| The above copyright notice and this permission notice shall be included in all | ||
| copies or substantial portions of the Software. | ||
|
|
||
| THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR | ||
| IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, | ||
| FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE | ||
| AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER | ||
| LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, | ||
| OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE | ||
| SOFTWARE. | ||
| */ | ||
|
|
||
| #ifndef HIP_TENSOR_AUDIO_AUGMENTATIONS_HPP | ||
| #define HIP_TENSOR_AUDIO_AUGMENTATIONS_HPP | ||
|
|
||
| #include "kernel/down_mixing.hpp" | ||
|
|
||
| #endif // HIP_TENSOR_AUDIO_AUGMENTATIONS_HPP |
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,73 @@ | ||
| #include <hip/hip_runtime.h> | ||
| #include "rpp_hip_common.hpp" | ||
|
|
||
| __global__ void down_mixing_hip_tensor(float *srcPtr, | ||
| uint srcStride, | ||
| float *dstPtr, | ||
| uint dstStride, | ||
| int2 *srcDimsTensor) | ||
|
|
||
| { | ||
| int id_x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; | ||
| int id_z = hipBlockIdx_z * hipBlockDim_z + hipThreadIdx_z; | ||
| int srcLength = srcDimsTensor[id_z].x; | ||
| int channels = srcDimsTensor[id_z].y; | ||
|
|
||
| if (id_x >= srcLength) | ||
| return; | ||
|
|
||
| float nomalizedWeight = 1.f / channels; | ||
| float outVal = 0.0f; | ||
| uint srcIdx = id_z * srcStride + id_x * channels; | ||
| int i = 0; | ||
| int alignedChannels = (channels / 8) * 8; | ||
| // if number of channels is a multiple of 8, do 8 pixel load till alignedChannels value | ||
| if (alignedChannels) | ||
| { | ||
| d_float8 outVal_f8; | ||
| outVal_f8.f4[0] = static_cast<float4>(0.0f); | ||
| outVal_f8.f4[1] = outVal_f8.f4[0]; | ||
| float4 normalizedWeight_f4 = static_cast<float4>(nomalizedWeight); | ||
| for(; i < alignedChannels; i += 8, srcIdx += 8) | ||
| { | ||
| d_float8 src_f8; | ||
| rpp_hip_load8_and_unpack_to_float8(srcPtr + srcIdx, &src_f8); | ||
| rpp_hip_math_multiply8_const(&src_f8, &src_f8, normalizedWeight_f4); | ||
|
There was a problem hiding this comment. @sampath1117 normalizedWeight is a constant which is just 1/channels. You don't need this multiply channels times in the loop. |
||
| rpp_hip_math_add8(&outVal_f8, &src_f8, &outVal_f8); | ||
| } | ||
| outVal_f8.f4[0] += outVal_f8.f4[1]; | ||
| outVal += (outVal_f8.f1[0] + outVal_f8.f1[1] + outVal_f8.f1[2] + outVal_f8.f1[3]); | ||
| } | ||
| // process remaining channels | ||
| for(; i < channels; i++, srcIdx++) | ||
| outVal += srcPtr[srcIdx] * nomalizedWeight; | ||
|
|
||
| uint dstIdx = id_z * dstStride + id_x; | ||
| dstPtr[dstIdx] = outVal; | ||
|
There was a problem hiding this comment. What is the range of values channels we can have? Any ways to vectorize the write here? There was a problem hiding this comment. I see that the RPPT_MAX_AUDIO_CHANNELS = 16 in the Resample PR. If thats the case, possibly running the 1 liner un-vectorized loop on L43 to do a "outVal += srcPtr[srcIdx]" should be fine followed by doing a "outVal *= nomalizedWeight" at the end.
There was a problem hiding this comment. Tried this experiment and it is giving a good 30% improvement in performance Previously Currently we have only 2 channel input in test suite and it goes to non vectorized loop, so it is loading 16 elements per thread and giving a 30% boost compared to unvectorized writes case There was a problem hiding this comment. I think this performance improvement was an anomoly Reran multiple times today and confirmed there was a degradation in performance |
||
| } | ||
|
|
||
| RppStatus hip_exec_down_mixing_tensor(Rpp32f *srcPtr, | ||
| RpptDescPtr srcDescPtr, | ||
| Rpp32f *dstPtr, | ||
| RpptDescPtr dstDescPtr, | ||
| Rpp32s *srcDimsTensor, | ||
| bool normalizeWeights, | ||
| rpp::Handle& handle) | ||
| { | ||
| Rpp32s globalThreads_x = dstDescPtr->strides.nStride; | ||
| Rpp32s globalThreads_y = 1; | ||
| Rpp32s globalThreads_z = dstDescPtr->n; | ||
|
|
||
| hipLaunchKernelGGL(down_mixing_hip_tensor, | ||
| dim3(ceil((Rpp32f)globalThreads_x/LOCAL_THREADS_X_1DIM), ceil((Rpp32f)globalThreads_y/LOCAL_THREADS_Y_1DIM), ceil((Rpp32f)globalThreads_z/LOCAL_THREADS_Z_1DIM)), | ||
| dim3(LOCAL_THREADS_X_1DIM, LOCAL_THREADS_Y_1DIM, LOCAL_THREADS_Z_1DIM), | ||
| 0, | ||
| handle.GetStream(), | ||
| srcPtr, | ||
| srcDescPtr->strides.nStride, | ||
| dstPtr, | ||
| dstDescPtr->strides.nStride, | ||
| reinterpret_cast<int2 *>(srcDimsTensor)); | ||
|
|
||
| return RPP_SUCCESS; | ||
| } | ||
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Why is this not vectorizeed?
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
we can only vectorize the channel code as how it is done in host code
Initially felt it might not be efficient to vectorize the channel loop in HIP
But explored further and vectorized the code now and checked the performance. It looks fine
Could not see improvement in performance though with current inputs
since we are testing with 2 channel input and vectorization only works with at least 8 channel input