Skip to content
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.

Sign up for GitHub

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

Jump to bottom

RPP Color Temperature on HOST and HIP #271

Merged
merged 20 commits into from
Feb 9, 2024
Merged

RPP Color Temperature on HOST and HIP #271

Show file tree
Hide file tree
Changes from 5 commits
Commits
Show all changes
20 commits
Select commit Hold shift + click to select a range
ad4e4a6
Initial commit - Color Temperature HOST Tensor
snehaa8 Dec 11, 2023
1ff57b9
Initial commit - Color Temperature HIP Tensor
snehaa8 Dec 11, 2023
13004e0
Add color temperature golden outputs
snehaa8 Dec 11, 2023
3bb3f3b
address review comments
snehaa8 Dec 13, 2023
11ea828
Merge pull request #207 from snehaa8/sn/color_temp_tensor
r-abishek Dec 13, 2023
15022a8
Merge branch 'master' into ar/opt_color_temperature
r-abishek Jan 3, 2024
c99321f
Use reinterpret_cast instead of static_cast
snehaa8 Jan 22, 2024
e4123c4
Combine templated functions to support all datatypes into one
snehaa8 Jan 22, 2024
d9b269d
Fix i8 datatype
snehaa8 Jan 23, 2024
e3983a0
Cleanup
snehaa8 Jan 23, 2024
a7f17f7
RPP Test Suite Upgrade 4 - CSV to BIN conversions for file size reduc…
r-abishek Jan 26, 2024
3019c8b
Merge branch 'develop' of https://github.com/ROCm/rpp into sn/color_t…
snehaa8 Jan 29, 2024
c700bf4
Fix PLN3 variant outputs
snehaa8 Jan 29, 2024
88ab397
Merge branch 'ar/opt_color_temperature' into sn/color_temp_tensor
snehaa8 Jan 29, 2024
9944e02
Merge branch 'develop' of https://github.com/GPUOpen-ProfessionalComp…
r-abishek Jan 29, 2024
45be1f8
Merge pull request #224 from snehaa8/sn/color_temp_tensor
r-abishek Jan 29, 2024
2a8e461
Update color_temperature.hpp license
r-abishek Jan 31, 2024
3e1be05
Delete color_temperature_u8_Tensor_PKD3.csv
r-abishek Jan 31, 2024
567233a
Delete color_temperature_u8_Tensor_PLN3.csv
r-abishek Jan 31, 2024
2a3b69c
Merge branch 'master' of https://github.com/GPUOpen-ProfessionalCompu…
r-abishek Feb 1, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 3 additions & 1 deletion include/rppdefs.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -113,7 +113,9 @@ typedef enum
/*! \brief Not enough memory \ingroup group_rppdefs */
RPP_ERROR_NOT_ENOUGH_MEMORY = -16,
/*! \brief Out of bound source ROI \ingroup group_rppdefs */
RPP_ERROR_OUT_OF_BOUND_SRC_ROI = -17
RPP_ERROR_OUT_OF_BOUND_SRC_ROI = -17,
/*! \brief Number of channels is invalid. (Needs to adhere to function specification.) \ingroup group_rppdefs */
RPP_ERROR_INVALID_CHANNELS = -18
} RppStatus;

/*! \brief RPP rppStatus_t type enums
Expand Down
42 changes: 42 additions & 0 deletions include/rppt_tensor_color_augmentations.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -415,6 +415,48 @@ RppStatus rppt_lut_host(RppPtr_t srcPtr, RpptDescPtr srcDescPtr, RppPtr_t dstPtr
RppStatus rppt_lut_gpu(RppPtr_t srcPtr, RpptDescPtr srcDescPtr, RppPtr_t dstPtr, RpptDescPtr dstDescPtr, RppPtr_t lutPtr, RpptROIPtr roiTensorPtrSrc, RpptRoiType roiType, rppHandle_t rppHandle);
#endif // GPU_SUPPORT

/*! \brief Color Temperature augmentation on HOST backend for a NCHW/NHWC layout tensor
* \details The color temperature augmentation does a image temperature adjustment operation, taking a pixel adjustment value as argument for each image in a batch of RGB(3 channel) with an NHWC/NCHW tensor layout.<br>
* - srcPtr depth ranges - Rpp8u (0 to 255), Rpp16f (0 to 1), Rpp32f (0 to 1), Rpp8s (-128 to 127).
* - dstPtr depth ranges - Will be same depth as srcPtr.
* \image html img150x150.jpg Sample Input
* \image html color_augmentations_color_temperature_img150x150.jpg Sample Output
* \param [in] srcPtr source tensor in HOST memory
* \param [in] srcDescPtr source tensor descriptor (Restrictions - numDims = 4, offsetInBytes >= 0, dataType = U8/F16/F32/I8, layout = NCHW/NHWC, c = 1/3)
* \param [out] dstPtr destination tensor in HOST memory
* \param [in] dstDescPtr destination tensor descriptor (Restrictions - numDims = 4, offsetInBytes >= 0, dataType = U8/F16/F32/I8, layout = NCHW/NHWC, c = same as that of srcDescPtr)
* \param [in] adjustmentValueTensor adjustment values for color temperature calculation (1D tensor of size sizeof(Rpp8s) * batchSize with -100 <= adjustmentValueTensor[i] >= 100 for each image in batch)
* \param [in] roiTensorSrc ROI data in HOST memory, for each image in source tensor (2D tensor of size batchSize * 4, in either format - XYWH(xy.x, xy.y, roiWidth, roiHeight) or LTRB(lt.x, lt.y, rb.x, rb.y))
* \param [in] roiType ROI type used (RpptRoiType::XYWH or RpptRoiType::LTRB)
* \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_color_temperature_host(RppPtr_t srcPtr, RpptDescPtr srcDescPtr, RppPtr_t dstPtr, RpptDescPtr dstDescPtr, Rpp8s *adjustmentValueTensor, RpptROIPtr roiTensorPtrSrc, RpptRoiType roiType, rppHandle_t rppHandle);

#ifdef GPU_SUPPORT
/*! \brief Color Temperature augmentation on HIP backend for a NCHW/NHWC layout tensor
* \details The color temperature augmentation does a image temperature adjustment operation, taking a pixel adjustment value as argument for each image in a batch of RGB(3 channel) with an NHWC/NCHW tensor layout.<br>
* - srcPtr depth ranges - Rpp8u (0 to 255), Rpp16f (0 to 1), Rpp32f (0 to 1), Rpp8s (-128 to 127).
* - dstPtr depth ranges - Will be same depth as srcPtr.
* \image html img150x150.jpg Sample Input
* \image html color_augmentations_color_temperature_img150x150.jpg Sample Output
* \param [in] srcPtr source tensor in HIP memory
* \param [in] srcDescPtr source tensor descriptor (Restrictions - numDims = 4, offsetInBytes >= 0, dataType = U8/F16/F32/I8, layout = NCHW/NHWC, c = 1/3)
* \param [out] dstPtr destination tensor in HIP memory
* \param [in] dstDescPtr destination tensor descriptor (Restrictions - numDims = 4, offsetInBytes >= 0, dataType = U8/F16/F32/I8, layout = NCHW/NHWC, c = same as that of srcDescPtr)
* \param [in] adjustmentValueTensor adjustment values for color temperature calculation (1D tensor of size sizeof(Rpp8s) * batchSize with -100 <= adjustmentValueTensor[i] >= 100 for each image in batch)
* \param [in] roiTensorSrc ROI data in HIP memory, for each image in source tensor (2D tensor of size batchSize * 4, in either format - XYWH(xy.x, xy.y, roiWidth, roiHeight) or LTRB(lt.x, lt.y, rb.x, rb.y))
* \param [in] roiType ROI type used (RpptRoiType::XYWH or RpptRoiType::LTRB)
* \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_color_temperature_gpu(RppPtr_t srcPtr, RpptDescPtr srcDescPtr, RppPtr_t dstPtr, RpptDescPtr dstDescPtr, Rpp32s *adjustmentValueTensor, RpptROIPtr roiTensorPtrSrc, RpptRoiType roiType, rppHandle_t rppHandle);
#endif // GPU_SUPPORT

/*! @}
*/

Expand Down
16 changes: 16 additions & 0 deletions src/include/cpu/rpp_cpu_common.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3028,6 +3028,22 @@ inline void compute_color_cast_12_host(__m128 *p, __m128 pMul, __m128 *pAdd)
p[2] = _mm_fmadd_ps(_mm_sub_ps(p[2], pAdd[2]), pMul, pAdd[2]); // color_cast adjustment Rs
}

inline void compute_color_temperature_48_host(__m256 *p, __m256 pAdj)
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

please rename the variable name to snake_case

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Rajy will get back on this.

Copy link
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Rajy, this is actually consistent with all other API. Camel case for variables and snake case for RPP API.

{
p[0] = _mm256_add_ps(p[0], pAdj); // color_temperature adjustment Rs
p[1] = _mm256_add_ps(p[1], pAdj); // color_temperature adjustment Rs
// no color_temperature adjustment Gs
p[4] = _mm256_sub_ps(p[4], pAdj); // color_temperature adjustment Bs
p[5] = _mm256_sub_ps(p[5], pAdj); // color_temperature adjustment Bs
}

inline void compute_color_temperature_24_host(__m256 *p, __m256 pAdj)
{
p[0] = _mm256_add_ps(p[0], pAdj); // color_temperature adjustment Rs
// no color_temperature adjustment Gs
p[2] = _mm256_sub_ps(p[2], pAdj); // color_temperature adjustment Bs
}

inline void compute_xywh_from_ltrb_host(RpptROIPtr roiPtrInput, RpptROIPtr roiPtrImage)
{
roiPtrImage->xywhROI.xy.x = roiPtrInput->ltrbROI.lt.x;
Expand Down
1 change: 1 addition & 0 deletions src/modules/cpu/host_tensor_color_augmentations.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -32,5 +32,6 @@ THE SOFTWARE.
#include "kernel/exposure.hpp"
#include "kernel/contrast.hpp"
#include "kernel/lut.hpp"
#include "kernel/color_temperature.hpp"

#endif // HOST_TENSOR_COLOR_AUGMENTATIONS_HPP
Loading