Audio HIP PR5 - Preemphasis Filter HIP Support

include/rppt_tensor_audio_augmentations.h

@@ -97,6 +97,24 @@ RppStatus rppt_to_decibels_host(RppPtr_t srcPtr, RpptDescPtr srcDescPtr, RppPtr_
  */
 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
+/*! \brief Pre Emphasis Filter augmentation on HIP backend
+* \details Pre Emphasis Filter augmentation for audio data
+* \param [in] srcPtr source tensor in HIP memory
+* \param [in] srcDescPtr source tensor descriptor (Restrictions - numDims = 3, offsetInBytes >= 0, dataType = F32)
+* \param [out] dstPtr destination tensor in HIP memory
+* \param [in] dstDescPtr destination tensor descriptor (Restrictions - numDims = 2, offsetInBytes >= 0, dataType = F32)
+* \param [in] srcLengthTensor source audio buffer length (1D tensor in HIP memory, of size batchSize)
+* \param [in] coeffTensor preemphasis coefficient (1D tensor in Pinned / HIP memory, of size batchSize)
+* \param [in] borderType border value policy
+* \param [in] rppHandle RPP HIP handle created with <tt>\ref rppCreateWithStreamAndBatchSize()</tt>
+* \return A <tt> \ref RppStatus</tt> enumeration.
+* \retval RPP_SUCCESS Successful completion.
+* \retval RPP_ERROR* Unsuccessful completion.
+*/
+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
 * \param[in] srcPtr source tensor in HOST memory
@@ -126,16 +144,12 @@ RppStatus rppt_down_mixing_host(RppPtr_t srcPtr, RpptDescPtr srcDescPtr, RppPtr_
  * \param [in] power exponent of the magnitude of the spectrum
  * \param [in] windowLength window size in number of samples
  * \param [in] windowStep step between the STFT windows in number of samples
- * \param [in] rppHandle RPP HOST handle created with <tt>\ref rppCreateWithBatchSize()</tt>
- * \return A <tt> \ref RppStatus</tt> enumeration.
- * \retval RPP_SUCCESS Successful completion.
- * \retval RPP_ERROR* Unsuccessful completion.
  */
 RppStatus rppt_spectrogram_host(RppPtr_t srcPtr, RpptDescPtr srcDescPtr, RppPtr_t dstPtr, RpptDescPtr dstDescPtr, Rpp32s *srcLengthTensor, bool centerWindows, bool reflectPadding, Rpp32f *windowFunction, Rpp32s nfft, Rpp32s power, Rpp32s windowLength, Rpp32s windowStep, rppHandle_t rppHandle);
 
 /*! \brief Mel filter bank augmentation HOST backend
  * \details Mel filter bank augmentation for audio data
- * \param[in] srcPtr source tensor in HOST memory
+ * \param[out] dstPtr srcPtr source tensor in HOST memory
  * \param[in] srcDescPtr source tensor descriptor (Restrictions - numDims = 3, offsetInBytes >= 0, dataType = F32, layout - NFT)
  * \param[out] dstPtr destination tensor in HOST memory
  * \param[in] dstDescPtr destination tensor descriptor (Restrictions - numDims = 3, offsetInBytes >= 0, dataType = F32, layout - NFT)

src/modules/hip/hip_tensor_audio_augmentations.hpp

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

src/modules/hip/kernel/pre_emphasis_filter.hpp

@@ -0,0 +1,74 @@
+#include <hip/hip_runtime.h>
+#include "rpp_hip_common.hpp"
+
+__device__ __forceinline__ 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,
+                                           uint srcStride,
+                                           float *dstPtr,
+                                           uint dstStride,
+                                           float *coeffTensor,
+                                           int *srcLengthTensor,
+                                           RpptAudioBorderType borderType)
+{
+    int id_x = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x) * 8;
+    int id_z = hipBlockIdx_z * hipBlockDim_z + hipThreadIdx_z;
+
+    if (id_x >= srcLengthTensor[id_z])
+        return;
+
+    uint srcIdx = (id_z * srcStride) + id_x;
+    uint dstIdx = (id_z * dstStride) + id_x;
+    float coeff = coeffTensor[id_z];
+
+    d_float8 src1_f8, src2_f8, dst_f8;
+    float4 coeff_f4 =  static_cast<float4>(coeff);
+
+    rpp_hip_load8_and_unpack_to_float8(srcPtr + srcIdx, &src1_f8);
+
+    if(id_x == 0)
+    {
+        src2_f8.f1[0] = (borderType == RpptAudioBorderType::CLAMP) ? src1_f8.f1[0] :
+                        (borderType == RpptAudioBorderType::REFLECT) ? src1_f8.f1[1] : 0.0f;
+        for(int i = 1; i < 8; i++)
+            src2_f8.f1[i] = src1_f8.f1[i - 1];
+    }
+    else
+        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,
+                                              Rpp32f *coeffTensor,
+                                              Rpp32s *srcLengthTensor,
+                                              RpptAudioBorderType borderType,
+                                              rpp::Handle& handle)
+{
+    Rpp32s globalThreads_x = (dstDescPtr->strides.nStride + 7) >> 3;
+    Rpp32s globalThreads_y = 1;
+    Rpp32s globalThreads_z = dstDescPtr->n;
+
+    hipLaunchKernelGGL(pre_emphasis_filter_tensor,
+                       dim3(ceil((float)globalThreads_x/LOCAL_THREADS_X_1DIM), ceil((float)globalThreads_y/LOCAL_THREADS_Y_1DIM), ceil((float)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,
+                       coeffTensor,
+                       srcLengthTensor,
+                       borderType);
+
+    return RPP_SUCCESS;
+}

src/modules/rppt_tensor_audio_augmentations.cpp

@@ -29,6 +29,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,
@@ -271,4 +275,45 @@ RppStatus rppt_resample_host(RppPtr_t srcPtr,
     }
 }
 
+/********************************************************************************************************************/
+/*********************************************** 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,
+                                       Rpp32s *srcLengthTensor,
+                                       Rpp32f *coeffTensor,
+                                       RpptAudioBorderType borderType,
+                                       rppHandle_t rppHandle)
+{
+#ifdef HIP_COMPILE
+
+    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,
+                                            coeffTensor,
+                                            srcLengthTensor,
+                                            borderType,
+                                            rpp::deref(rppHandle));
+    }
+    else
+    {
+        return RPP_ERROR_NOT_IMPLEMENTED;
+    }
+
+    return RPP_SUCCESS;
+#elif defined(OCL_COMPILE)
+    return RPP_ERROR_NOT_IMPLEMENTED;
+#endif // backend
+}
+
 #endif // AUDIO_SUPPORT

utilities/test_suite/HIP/CMakeLists.txt

@@ -58,6 +58,10 @@ find_package(hip QUIET)
 find_package(OpenCV QUIET)
 find_package(TurboJpeg QUIET)
 find_package(NIFTI QUIET)
+find_library(libsnd_LIBS
+    NAMES sndfile libsndfile
+    PATHS ${CMAKE_SYSTEM_PREFIX_PATH} ${LIBSND_ROOT_DIR} "/usr/local"
+    PATH_SUFFIXES lib lib64)
 
 # OpenMP
 find_package(OpenMP REQUIRED)
@@ -111,4 +115,23 @@ if(NIFTI_FOUND AND OpenCV_FOUND)
     target_link_libraries(Tensor_voxel_hip ${OpenCV_LIBS} -lturbojpeg -lrpp ${hip_LIBRARIES} pthread ${LINK_LIBRARY_LIST} hip::device ${NIFTI_PACKAGE_PREFIX}NIFTI::${NIFTI_PACKAGE_PREFIX}niftiio)
 else()
     message("-- ${Yellow}Warning: libniftiio must be installed to install ${PROJECT_NAME}/Tensor_voxel_hip successfully!${ColourReset}")
-endif()
+endif()
+
+if(NOT libsnd_LIBS)
+    message("-- ${Yellow}Warning: libsndfile must be installed to install ${PROJECT_NAME}/Tensor_audio_host successfully!${ColourReset}")
+else()
+    message("-- ${Green}${PROJECT_NAME} set to build with rpp and libsndfile ${ColourReset}")
+    set(COMPILER_FOR_HIP ${ROCM_PATH}/bin/hipcc)
+    set(CMAKE_CXX_COMPILER ${COMPILER_FOR_HIP})
+    include_directories(${ROCM_PATH}/include ${ROCM_PATH}/include/rpp /usr/local/include)
+    link_directories(${ROCM_PATH}/lib /usr/local/lib)
+    include_directories(${SndFile_INCLUDE_DIRS})
+    link_directories(${SndFile_LIBRARIES_DIR} /usr/local/lib/)
+
+    add_executable(Tensor_audio_hip Tensor_audio_hip.cpp)
+    set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=gnu++17")
+    if(NOT APPLE)
+        set(LINK_LIBRARY_LIST ${LINK_LIBRARY_LIST} stdc++fs)
+    endif()
+    target_link_libraries(Tensor_audio_hip ${libsnd_LIBS} -lsndfile -lrpp pthread ${LINK_LIBRARY_LIST})
+endif()