-
Notifications
You must be signed in to change notification settings - Fork 39
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
RPP Audio Support HIP - Downmixing #401
Merged
kiritigowda
merged 67 commits into
ROCm:develop
from
r-abishek:ar/audio_support_4_down_mixing_hip
Aug 2, 2024
Merged
Changes from 62 commits
Commits
Show all changes
67 commits
Select commit
Hold shift + click to select a range
c33af22
Bump rocm-docs-core[api_reference] from 0.35.0 to 0.35.1 in /docs/sph…
dependabot[bot] 14f6334
Bump rocm-docs-core[api_reference] from 0.35.1 to 0.36.0 in /docs/sph…
dependabot[bot] 95c3272
Merge branch 'master' into develop
kiritigowda 3973c34
added api support for ToDecibels HIP kernel
sampath1117 3f08f90
added test suite support for audio in HIP
sampath1117 b9c0788
added profiler support for hip test suite
sampath1117 641f653
Docs - Bump rocm-docs-core[api_reference] from 0.36.0 to 0.37.0 in /d…
dependabot[bot] 5568573
Link cleanup (#326)
LisaDelaney a6749ba
Update notes
LisaDelaney a255906
Docs - Bump rocm-docs-core[api_reference] from 0.37.0 to 0.37.1 in /d…
dependabot[bot] d3df761
RPP Voxel Flip on HIP and HOST (#285)
r-abishek ebecb42
RPP Vignette Tensor on HOST and HIP (#311)
r-abishek fc1410b
Bump rocm-docs-core[api_reference] from 0.37.1 to 0.38.0 in /docs/sph…
dependabot[bot] cb3d539
added initial api and test suite support for downmixing hip kernel
sampath1117 7fe2a0e
Merge branch 'develop' into sr/downmixing_hip
sampath1117 e80533a
initial working commit for downmixing hip kernel
sampath1117 50a743f
added support fo copy input to output when channels is 1
sampath1117 fe06106
minor code cleanup
sampath1117 3ebd7c3
RPP Tensor Audio Support - Resample (#310)
r-abishek 76f31df
Docs - Missing input and output images for Doxygen (#331)
r-abishek b83f910
Scratch buffers rename for HOST and HIP (#324)
r-abishek ebeb131
Update CMakeLists.txt
kiritigowda 7c194b2
Merge branch 'develop' into sr/downmixing_hip
sampath1117 cb06f7f
added missing hipDeviceSynchronize() in test suite
sampath1117 0e51993
removed f16 includes since not needed for audio
sampath1117 004e1d6
restructured python test suite
sampath1117 b8f5c60
fixed spacing in Doxygen
sampath1117 1147bfe
Update CMakeLists.txt
kiritigowda 9d48447
Merge remote-tracking branch 'develop' into sr/downmixing_hip
sampath1117 5e3fc7a
Bump rocm-docs-core[api_reference] from 0.38.1 to 1.0.0 in /docs/sphi…
dependabot[bot] b6b7cc5
Bump rocm-docs-core[api_reference] from 1.0.0 to 1.1.0 in /docs/sphin…
dependabot[bot] e16ad7a
RPP Gaussian Noise Voxel Tensor on HOST and HIP (#323)
r-abishek 9394c78
Merge branch 'develop' into sr/downmixing_hip
sampath1117 06263a5
modify CHECK to CHECK_RETURN_STATUS
sampath1117 a7e71a3
Merge branch 'develop' into sr/downmixing_hip
sampath1117 77e14ef
Minor common-fixes for HIP (#345)
r-abishek 34f3f6d
Readme Updates: --usecase=rocm (#349)
kiritigowda ab52683
RPP Tensor Audio Support - Spectrogram (#312)
r-abishek ee0d6fe
Update CHANGELOG.md (#352)
r-abishek 2decd32
RPP Tensor Audio Support - Slice (#325)
r-abishek 30ce1d6
RPP Tensor Audio Support - MelFilterBank (#332)
r-abishek 64ae74f
RPP Tensor Normalize ND on HOST and HIP (#335)
r-abishek 1a3015c
SWDEV-459739 - Remove the package obsolete setting (#353)
raramakr b926816
Merge branch 'develop' into sr/downmixing_hip
sampath1117 7cb3c03
changed globalThreads_z to use batchsize from description pointer
sampath1117 4cb8d4b
Audio support merge commit fixes (#354)
r-abishek 8aab10c
Merge branch 'develop' into sr/downmixing_hip
sampath1117 9d68c49
removed if else block based on channels inside kernel
sampath1117 195e4a4
rename instances of tensor_hip_audio to tensor_audio_hip
sampath1117 807bbe8
modified verify_output to have different cutoff for HIP and HOST back…
sampath1117 b8264f7
vectorized channels loop in hip kernel
sampath1117 5cf7360
Merge branch 'develop' into sr/downmixing_hip
sampath1117 d62e9cb
Merge branch 'develop' into sr/downmixing_hip
sampath1117 239bf76
removed the multiplication with normalizeWeight for every channel value
sampath1117 fd21921
vectorized the writes for hip kernel
sampath1117 23072da
moved constant compute outside the loop
sampath1117 3b00c64
reverted back to unvectorized writes version
sampath1117 0c71365
Merge branch 'develop' into sr/downmixing_hip
sampath1117 a2e321e
Merge pull request #261 from sampath1117/sr/downmixing_hip
r-abishek 315b64b
audio test suite changes for python 2 compatibility
sampath1117 7f9635a
added validation checks for numDims
sampath1117 611fb3f
Merge pull request #295 from sampath1117/sr/down_mixing_hip_pr_changes
r-abishek 76ee4fe
Merge branch 'develop' into ar/audio_support_4_down_mixing_hip
kiritigowda 0745bc0
Merge branch 'develop' into sr/down_mixing_develop_merge
sampath1117 4aabed4
removed duplicate code added in merge
sampath1117 940e7a0
Merge pull request #298 from sampath1117/sr/down_mixing_develop_merge
r-abishek af1a238
Merge branch 'develop' of github.com:ROCm/rpp into ar/audio_support_4…
r-abishek File filter
Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
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 |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,72 @@ | ||
#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 outVal = 0.0f; | ||
uint srcIdx = id_z * srcStride + id_x * channels; | ||
int i = 0; | ||
int alignedChannels = (channels / 8) * 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]; | ||
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_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]; | ||
outVal *= (1.f / channels); | ||
|
||
uint dstIdx = id_z * dstStride + id_x; | ||
dstPtr[dstIdx] = outVal; | ||
} | ||
|
||
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; | ||
} |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Oops, something went wrong.
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
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.
This is rounding down. Is that what you want?
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.
Yes Rajy
Since we want to do 8 pixel vectorized processing for alignedChannels value