Audio HIP PR5 - Preemphasis Filter HIP Support

.readthedocs.yaml

@@ -15,4 +15,4 @@ python:
 build:
    os: ubuntu-22.04
    tools:
-      python: "3.8"
+      python: "3.10"

CMakeLists.txt

@@ -27,7 +27,7 @@ set(CMAKE_CXX_COMPILER clang++)
 set(CMAKE_CXX_STANDARD 17)
 
 # RPP Version
-set(VERSION "1.6.0")
+set(VERSION "1.7.0")
 
 # Set Project Version and Language
 project(rpp VERSION ${VERSION} LANGUAGES CXX)

docs/sphinx/requirements.in

@@ -1 +1 @@
-rocm-docs-core[api_reference]==0.38.1
+rocm-docs-core[api_reference]==1.0.0

docs/sphinx/requirements.txt

@@ -110,7 +110,7 @@ requests==2.28.2
     # via
     #   pygithub
     #   sphinx
-rocm-docs-core[api-reference]==0.38.1
+rocm-docs-core[api-reference]==1.0.0
     # via -r requirements.in
 smmap==5.0.0
     # via gitdb

include/rppt_tensor_audio_augmentations.h

@@ -94,6 +94,9 @@ RppStatus rppt_to_decibels_host(RppPtr_t srcPtr, RpptDescPtr srcDescPtr, RppPtr_
  * \retval RPP_ERROR* Unsuccessful completion.
  */
 RppStatus rppt_pre_emphasis_filter_host(RppPtr_t srcPtr, RpptDescPtr srcDescPtr, RppPtr_t dstPtr, RpptDescPtr dstDescPtr, Rpp32s *srcLengthTensor, Rpp32f *coeffTensor, RpptAudioBorderType borderType, rppHandle_t rppHandle);
+#ifdef GPU_SUPPORT
+RppStatus rppt_pre_emphasis_filter_gpu(RppPtr_t srcPtr, RpptDescPtr srcDescPtr, RppPtr_t dstPtr, RpptDescPtr dstDescPtr, Rpp32s *srcLengthTensor, Rpp32f *coeffTensor, RpptAudioBorderType borderType, rppHandle_t rppHandle);
+#endif // GPU_SUPPORT
 
 /*! \brief Down Mixing augmentation on HOST backend
 * \details Down Mixing augmentation for audio data

src/modules/hip/hip_tensor_audio_augmentations.hpp

@@ -0,0 +1,6 @@
+#ifndef HIP_TENSOR_AUDIO_HPP
+#define HIP_TENSOR_AUDIO_HPP
+
+#include "kernel/pre_emphasis_filter.hpp"
+
+#endif // HIP_TENSOR_AUDIO_HPP

src/modules/hip/kernel/pre_emphasis_filter.hpp

@@ -0,0 +1,78 @@
+#include <hip/hip_runtime.h>
+#include "rpp_hip_common.hpp"
+
+__device__ void pre_emphasis_filter_hip_compute(d_float8 *src1_f8, d_float8 *src2_f8, d_float8 *dst_f8, float4 *coeff_f4)
+{
+    dst_f8->f4[0] = src1_f8->f4[0]  - *coeff_f4 * src2_f8->f4[0];
+    dst_f8->f4[1] = src1_f8->f4[1]  - *coeff_f4 * src2_f8->f4[1];
+}
+
+__global__ void pre_emphasis_filter_tensor(float *srcPtr,
+                                           uint2 srcStridesNH,
+                                           float *dstPtr,
+                                           uint2 dstStridesNH,
+                                           RpptImagePatchPtr srcDims,
+                                           float *coeffTensor,
+                                           RpptAudioBorderType borderType)
+{
+    int id_x = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x) * 8 + 1;
+    int id_y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
+    int id_z = hipBlockIdx_z * hipBlockDim_z + hipThreadIdx_z;
+
+    if ((id_x >= srcDims[id_z].width) || (id_y >= srcDims[id_z].height))
+    {
+        return;
+    }
+
+    uint srcIdx = (id_z * srcStridesNH.x) + (id_y * srcStridesNH.y) + id_x;
+    uint dstIdx = (id_z * dstStridesNH.x) + (id_y * dstStridesNH.y) + id_x;
+
+    float4 coeff_f4 = (float4)coeffTensor[id_z];
+    d_float8 src1_f8, src2_f8, dst_f8;
+    rpp_hip_load8_and_unpack_to_float8(srcPtr + srcIdx, &src1_f8);
+    rpp_hip_load8_and_unpack_to_float8(srcPtr + srcIdx - 1, &src2_f8);
+    pre_emphasis_filter_hip_compute(&src1_f8, &src2_f8, &dst_f8, &coeff_f4);
+    rpp_hip_pack_float8_and_store8(dstPtr + dstIdx, &dst_f8);
+}
+
+RppStatus hip_exec_pre_emphasis_filter_tensor(Rpp32f *srcPtr,
+                                              RpptDescPtr srcDescPtr,
+                                              Rpp32f *dstPtr,
+                                              RpptDescPtr dstDescPtr,
+                                              RpptImagePatchPtr srcDims,
+                                              RpptAudioBorderType borderType,
+                                              rpp::Handle& handle)
+{
+    int globalThreads_x = (dstDescPtr->w + 7) >> 3;
+    int globalThreads_y = dstDescPtr->h;
+    int globalThreads_z = handle.GetBatchSize();
+
+    float *coeff = handle.GetInitHandle()->mem.mgpu.floatArr[0].floatmem;
+
+    for(int i = 0; i < srcDescPtr->n; i++)
+    {
+        int id_x = i * srcDescPtr->strides.nStride;
+        if(borderType == RpptAudioBorderType::ZERO)
+            dstPtr[id_x] = srcPtr[id_x];
+        else
+        {
+            float border = (borderType == RpptAudioBorderType::CLAMP) ? srcPtr[id_x] : srcPtr[id_x + 1];
+            dstPtr[id_x] = srcPtr[id_x] - coeff[id_x] * border;
+        }
+    }
+
+    hipLaunchKernelGGL(pre_emphasis_filter_tensor,
+                       dim3(ceil((float)globalThreads_x/LOCAL_THREADS_X), ceil((float)globalThreads_y/LOCAL_THREADS_Y), ceil((float)globalThreads_z/LOCAL_THREADS_Z)),
+                       dim3(LOCAL_THREADS_X, LOCAL_THREADS_Y, LOCAL_THREADS_Z),
+                       0,
+                       handle.GetStream(),
+                       srcPtr,
+                       make_uint2(srcDescPtr->strides.nStride, srcDescPtr->strides.hStride),
+                       dstPtr,
+                       make_uint2(dstDescPtr->strides.nStride, dstDescPtr->strides.hStride),
+                       srcDims,
+                       coeff,
+                       borderType);
+
+    return RPP_SUCCESS;
+}

src/modules/rppt_tensor_audio_augmentations.cpp

@@ -27,6 +27,10 @@ SOFTWARE.
 #include "rppt_tensor_audio_augmentations.h"
 #include "cpu/host_tensor_audio_augmentations.hpp"
 
+#ifdef HIP_COMPILE
+    #include "hip/hip_tensor_audio_augmentations.hpp"
+#endif // HIP_COMPILE
+
 /******************** non_silent_region_detection ********************/
 
 RppStatus rppt_non_silent_region_detection_host(RppPtr_t srcPtr,
@@ -186,3 +190,42 @@ RppStatus rppt_resample_host(RppPtr_t srcPtr,
         return RPP_ERROR_NOT_IMPLEMENTED;
     }
 }
+
+/********************************************************************************************************************/
+/*********************************************** RPP_GPU_SUPPORT = ON ***********************************************/
+/********************************************************************************************************************/
+
+#ifdef GPU_SUPPORT
+
+/******************** pre_emphasis_filter ********************/
+
+RppStatus rppt_pre_emphasis_filter_gpu(RppPtr_t srcPtr,
+                                       RpptDescPtr srcDescPtr,
+                                       RppPtr_t dstPtr,
+                                       RpptDescPtr dstDescPtr,
+                                       RpptImagePatchPtr srcDims,
+                                       Rpp32f *coeffTensor,
+                                       RpptAudioBorderType borderType,
+                                       rppHandle_t rppHandle)
+{
+#ifdef HIP_COMPILE
+    Rpp32u paramIndex = 0;
+    copy_param_float(coeffTensor, rpp::deref(rppHandle), paramIndex++);
+
+    if ((srcDescPtr->dataType == RpptDataType::F32) && (dstDescPtr->dataType == RpptDataType::F32))
+    {
+        hip_exec_pre_emphasis_filter_tensor(static_cast<Rpp32f*>(srcPtr),
+                                            srcDescPtr,
+                                            static_cast<Rpp32f*>(dstPtr),
+                                            dstDescPtr,
+                                            srcDims,
+                                            borderType,
+                                            rpp::deref(rppHandle));
+    }
+    return RPP_SUCCESS;
+#elif defined(OCL_COMPILE)
+    return RPP_ERROR_NOT_IMPLEMENTED;
+#endif // backend
+}
+
+#endif // GPU_SUPPORT