From 5c423ab686153dcfeaaf4364743837382e3a27ba Mon Sep 17 00:00:00 2001 From: Abishek <52214183+r-abishek@users.noreply.github.com> Date: Fri, 9 Feb 2024 15:00:48 -0800 Subject: [PATCH] RPP Color Temperature on HOST and HIP (#271) * Initial commit - Color Temperature HOST Tensor * Initial commit - Color Temperature HIP Tensor * Add color temperature golden outputs * address review comments * Use reinterpret_cast instead of static_cast * Combine templated functions to support all datatypes into one (got minor perf difference of order 3%) Also fixes indentation * Fix i8 datatype * Cleanup * RPP Test Suite Upgrade 4 - CSV to BIN conversions for file size reduction (#293) * change golden outputs from .csv files to .bin files * Changed comparision funtions to use .bin files * Address review comments * minor change * Address review comments * minor change --------- Co-authored-by: HazarathKumarM * Fix PLN3 variant outputs Also modifies reference outputs * Update color_temperature.hpp license * Delete color_temperature_u8_Tensor_PKD3.csv * Delete color_temperature_u8_Tensor_PLN3.csv --------- Co-authored-by: snehaa8 Co-authored-by: HazarathKumarM Co-authored-by: Snehaa-Giridharan <118163708+snehaa8@users.noreply.github.com> --- include/rppdefs.h | 4 +- include/rppt_tensor_color_augmentations.h | 42 + src/include/cpu/rpp_cpu_common.hpp | 16 + .../cpu/host_tensor_color_augmentations.hpp | 1 + src/modules/cpu/kernel/color_temperature.hpp | 1035 +++++++++++++++++ .../hip/hip_tensor_color_augmentations.hpp | 1 + src/modules/hip/kernel/color_temperature.hpp | 223 ++++ .../rppt_tensor_color_augmentations.cpp | 139 ++- utilities/test_suite/HIP/Tensor_hip.cpp | 18 +- utilities/test_suite/HIP/runTests.py | 4 +- utilities/test_suite/HOST/Tensor_host.cpp | 19 +- utilities/test_suite/HOST/runTests.py | 4 +- .../color_temperature_u8_Tensor.bin | Bin 0 -> 205200 bytes utilities/test_suite/rpp_test_suite_common.h | 1 + 14 files changed, 1497 insertions(+), 10 deletions(-) create mode 100644 src/modules/cpu/kernel/color_temperature.hpp create mode 100644 src/modules/hip/kernel/color_temperature.hpp create mode 100644 utilities/test_suite/REFERENCE_OUTPUT/color_temperature/color_temperature_u8_Tensor.bin diff --git a/include/rppdefs.h b/include/rppdefs.h index 996c4e3c7..b0baf7d34 100644 --- a/include/rppdefs.h +++ b/include/rppdefs.h @@ -115,7 +115,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 diff --git a/include/rppt_tensor_color_augmentations.h b/include/rppt_tensor_color_augmentations.h index deabd885d..99909cb42 100644 --- a/include/rppt_tensor_color_augmentations.h +++ b/include/rppt_tensor_color_augmentations.h @@ -417,6 +417,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.
+ * - 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 \ref rppCreateWithBatchSize() + * \return A \ref RppStatus 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.
+ * - 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 \ref rppCreateWithStreamAndBatchSize() + * \return A \ref RppStatus 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 + /*! @} */ diff --git a/src/include/cpu/rpp_cpu_common.hpp b/src/include/cpu/rpp_cpu_common.hpp index 67c34de70..711cdbb5e 100644 --- a/src/include/cpu/rpp_cpu_common.hpp +++ b/src/include/cpu/rpp_cpu_common.hpp @@ -3032,6 +3032,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) +{ + 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; diff --git a/src/modules/cpu/host_tensor_color_augmentations.hpp b/src/modules/cpu/host_tensor_color_augmentations.hpp index 19e0b471c..aba3b8158 100644 --- a/src/modules/cpu/host_tensor_color_augmentations.hpp +++ b/src/modules/cpu/host_tensor_color_augmentations.hpp @@ -34,5 +34,6 @@ SOFTWARE. #include "kernel/exposure.hpp" #include "kernel/contrast.hpp" #include "kernel/lut.hpp" +#include "kernel/color_temperature.hpp" #endif // HOST_TENSOR_COLOR_AUGMENTATIONS_HPP diff --git a/src/modules/cpu/kernel/color_temperature.hpp b/src/modules/cpu/kernel/color_temperature.hpp new file mode 100644 index 000000000..1358ac800 --- /dev/null +++ b/src/modules/cpu/kernel/color_temperature.hpp @@ -0,0 +1,1035 @@ +/* +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. +*/ + +#include "rppdefs.h" +#include "rpp_cpu_simd.hpp" +#include "rpp_cpu_common.hpp" + +RppStatus color_temperature_u8_u8_host_tensor(Rpp8u *srcPtr, + RpptDescPtr srcDescPtr, + Rpp8u *dstPtr, + RpptDescPtr dstDescPtr, + Rpp8s *adjustmentValueTensor, + RpptROIPtr roiTensorPtrSrc, + RpptRoiType roiType, + RppLayoutParams layoutParams) +{ + RpptROI roiDefault = {0, 0, (Rpp32s)srcDescPtr->w, (Rpp32s)srcDescPtr->h}; + + omp_set_dynamic(0); +#pragma omp parallel for num_threads(dstDescPtr->n) + for(int batchCount = 0; batchCount < dstDescPtr->n; batchCount++) + { + RpptROI roi; + RpptROIPtr roiPtrInput = &roiTensorPtrSrc[batchCount]; + compute_roi_validation_host(roiPtrInput, &roi, &roiDefault, roiType); + + Rpp32f adjustmentValue = adjustmentValueTensor[batchCount]; + + Rpp8u *srcPtrImage, *dstPtrImage; + srcPtrImage = srcPtr + batchCount * srcDescPtr->strides.nStride; + dstPtrImage = dstPtr + batchCount * dstDescPtr->strides.nStride; + + Rpp32u bufferLength = roi.xywhROI.roiWidth * layoutParams.bufferMultiplier; + + __m256 pAdj = _mm256_set1_ps(adjustmentValue); + + Rpp8u *srcPtrChannel, *dstPtrChannel; + srcPtrChannel = srcPtrImage + (roi.xywhROI.xy.y * srcDescPtr->strides.hStride) + (roi.xywhROI.xy.x * layoutParams.bufferMultiplier); + dstPtrChannel = dstPtrImage; + + // Color Temperature with fused output-layout toggle (NHWC -> NCHW) + if ((srcDescPtr->c == 3) && (srcDescPtr->layout == RpptLayout::NHWC) && (dstDescPtr->layout == RpptLayout::NCHW)) + { + Rpp32u alignedLength = (bufferLength / 48) * 48; + + Rpp8u *srcPtrRow, *dstPtrRowR, *dstPtrRowG, *dstPtrRowB; + srcPtrRow = srcPtrChannel; + dstPtrRowR = dstPtrChannel; + dstPtrRowG = dstPtrRowR + dstDescPtr->strides.cStride; + dstPtrRowB = dstPtrRowG + dstDescPtr->strides.cStride; + + for(int i = 0; i < roi.xywhROI.roiHeight; i++) + { + Rpp8u *srcPtrTemp, *dstPtrTempR, *dstPtrTempG, *dstPtrTempB; + srcPtrTemp = srcPtrRow; + dstPtrTempR = dstPtrRowR; + dstPtrTempG = dstPtrRowG; + dstPtrTempB = dstPtrRowB; + + int vectorLoopCount = 0; + for (; vectorLoopCount < alignedLength; vectorLoopCount += 48) + { + __m256 p[6]; + + rpp_simd_load(rpp_load48_u8pkd3_to_f32pln3_avx, srcPtrTemp, p); // simd loads + compute_color_temperature_48_host(p, pAdj); // color_temperature adjustment + rpp_simd_store(rpp_store48_f32pln3_to_u8pln3_avx, dstPtrTempR, dstPtrTempG, dstPtrTempB, p); // simd stores + + srcPtrTemp += 48; + dstPtrTempR += 16; + dstPtrTempG += 16; + dstPtrTempB += 16; + } + for (; vectorLoopCount < bufferLength; vectorLoopCount += 3) + { + *dstPtrTempR++ = (Rpp8u) RPPPIXELCHECK(srcPtrTemp[0] + adjustmentValue); + *dstPtrTempG++ = (Rpp8u) RPPPIXELCHECK(srcPtrTemp[1]); + *dstPtrTempB++ = (Rpp8u) RPPPIXELCHECK(srcPtrTemp[2] - adjustmentValue); + + srcPtrTemp += 3; + } + + srcPtrRow += srcDescPtr->strides.hStride; + dstPtrRowR += dstDescPtr->strides.hStride; + dstPtrRowG += dstDescPtr->strides.hStride; + dstPtrRowB += dstDescPtr->strides.hStride; + } + } + + // Color Temperature with fused output-layout toggle (NCHW -> NHWC) + else if ((srcDescPtr->c == 3) && (srcDescPtr->layout == RpptLayout::NCHW) && (dstDescPtr->layout == RpptLayout::NHWC)) + { + Rpp32u alignedLength = (bufferLength / 48) * 48; + + Rpp8u *srcPtrRowR, *srcPtrRowG, *srcPtrRowB, *dstPtrRow; + srcPtrRowR = srcPtrChannel; + srcPtrRowG = srcPtrRowR + srcDescPtr->strides.cStride; + srcPtrRowB = srcPtrRowG + srcDescPtr->strides.cStride; + dstPtrRow = dstPtrChannel; + + for(int i = 0; i < roi.xywhROI.roiHeight; i++) + { + Rpp8u *srcPtrTempR, *srcPtrTempG, *srcPtrTempB, *dstPtrTemp; + srcPtrTempR = srcPtrRowR; + srcPtrTempG = srcPtrRowG; + srcPtrTempB = srcPtrRowB; + dstPtrTemp = dstPtrRow; + + int vectorLoopCount = 0; + for (; vectorLoopCount < alignedLength; vectorLoopCount += 16) + { + __m256 p[6]; + + rpp_simd_load(rpp_load48_u8pln3_to_f32pln3_avx, srcPtrTempR, srcPtrTempG, srcPtrTempB, p); // simd loads + compute_color_temperature_48_host(p, pAdj); // color_temperature adjustment + rpp_simd_store(rpp_store48_f32pln3_to_u8pkd3_avx, dstPtrTemp, p); // simd stores + + srcPtrTempR += 16; + srcPtrTempG += 16; + srcPtrTempB += 16; + dstPtrTemp += 48; + } + for (; vectorLoopCount < bufferLength; vectorLoopCount++) + { + dstPtrTemp[0] = (Rpp8u) RPPPIXELCHECK(*srcPtrTempR + adjustmentValue); + dstPtrTemp[1] = (Rpp8u) RPPPIXELCHECK(*srcPtrTempG); + dstPtrTemp[2] = (Rpp8u) RPPPIXELCHECK(*srcPtrTempB - adjustmentValue); + + dstPtrTemp += 3; + srcPtrTempR++; + srcPtrTempG++; + srcPtrTempB++; + } + + srcPtrRowR += srcDescPtr->strides.hStride; + srcPtrRowG += srcDescPtr->strides.hStride; + srcPtrRowB += srcDescPtr->strides.hStride; + dstPtrRow += dstDescPtr->strides.hStride; + } + } + + // Color Temperature with fused output-layout toggle (NHWC -> NHWC) + else if ((srcDescPtr->c == 3) && (srcDescPtr->layout == RpptLayout::NHWC) && (dstDescPtr->layout == RpptLayout::NHWC)) + { + Rpp32u alignedLength = (bufferLength / 48) * 48; + + Rpp8u *srcPtrRow, *dstPtrRow; + srcPtrRow = srcPtrChannel; + dstPtrRow = dstPtrChannel; + + for(int i = 0; i < roi.xywhROI.roiHeight; i++) + { + Rpp8u *srcPtrTemp, *dstPtrTemp; + srcPtrTemp = srcPtrRow; + dstPtrTemp = dstPtrRow; + + int vectorLoopCount = 0; + for (; vectorLoopCount < alignedLength; vectorLoopCount += 48) + { + __m256 p[6]; + + rpp_simd_load(rpp_load48_u8pkd3_to_f32pln3_avx, srcPtrTemp, p); // simd loads + compute_color_temperature_48_host(p, pAdj); // color_temperature adjustment + rpp_simd_store(rpp_store48_f32pln3_to_u8pkd3_avx, dstPtrTemp, p); // simd stores + + srcPtrTemp += 48; + dstPtrTemp += 48; + } + for (; vectorLoopCount < bufferLength; vectorLoopCount += 3) + { + dstPtrTemp[0] = (Rpp8u) RPPPIXELCHECK(srcPtrTemp[0] + adjustmentValue); + dstPtrTemp[1] = (Rpp8u) RPPPIXELCHECK(srcPtrTemp[1]); + dstPtrTemp[2] = (Rpp8u) RPPPIXELCHECK(srcPtrTemp[2] - adjustmentValue); + + srcPtrTemp += 3; + dstPtrTemp += 3; + } + + srcPtrRow += srcDescPtr->strides.hStride; + dstPtrRow += dstDescPtr->strides.hStride; + } + } + + // Color Temperature with fused output-layout toggle (NCHW -> NCHW) + else if ((srcDescPtr->c == 3) && (srcDescPtr->layout == RpptLayout::NCHW) && (dstDescPtr->layout == RpptLayout::NCHW)) + { + Rpp32u alignedLength = (bufferLength / 48) * 48; + + Rpp8u *srcPtrRowR, *srcPtrRowG, *srcPtrRowB, *dstPtrRowR, *dstPtrRowG, *dstPtrRowB; + srcPtrRowR = srcPtrChannel; + srcPtrRowG = srcPtrRowR + srcDescPtr->strides.cStride; + srcPtrRowB = srcPtrRowG + srcDescPtr->strides.cStride; + dstPtrRowR = dstPtrChannel; + dstPtrRowG = dstPtrRowR + dstDescPtr->strides.cStride; + dstPtrRowB = dstPtrRowG + dstDescPtr->strides.cStride; + + for(int i = 0; i < roi.xywhROI.roiHeight; i++) + { + Rpp8u *srcPtrTempR, *srcPtrTempG, *srcPtrTempB, *dstPtrTempR, *dstPtrTempG, *dstPtrTempB; + srcPtrTempR = srcPtrRowR; + srcPtrTempG = srcPtrRowG; + srcPtrTempB = srcPtrRowB; + dstPtrTempR = dstPtrRowR; + dstPtrTempG = dstPtrRowG; + dstPtrTempB = dstPtrRowB; + + int vectorLoopCount = 0; + for (; vectorLoopCount < alignedLength; vectorLoopCount += 16) + { + __m256 p[6]; + + rpp_simd_load(rpp_load48_u8pln3_to_f32pln3_avx, srcPtrTempR, srcPtrTempG, srcPtrTempB, p); // simd loads + compute_color_temperature_48_host(p, pAdj); // color_temperature adjustment + rpp_simd_store(rpp_store48_f32pln3_to_u8pln3_avx, dstPtrTempR, dstPtrTempG, dstPtrTempB, p); // simd stores + + srcPtrTempR += 16; + srcPtrTempG += 16; + srcPtrTempB += 16; + dstPtrTempR += 16; + dstPtrTempG += 16; + dstPtrTempB += 16; + } + for (; vectorLoopCount < bufferLength; vectorLoopCount++) + { + *dstPtrTempR++ = (Rpp8u) RPPPIXELCHECK(*srcPtrTempR + adjustmentValue); + *dstPtrTempG++ = (Rpp8u) RPPPIXELCHECK(*srcPtrTempG); + *dstPtrTempB++ = (Rpp8u) RPPPIXELCHECK(*srcPtrTempB - adjustmentValue); + + srcPtrTempR++; + srcPtrTempG++; + srcPtrTempB++; + } + + srcPtrRowR += srcDescPtr->strides.hStride; + srcPtrRowG += srcDescPtr->strides.hStride; + srcPtrRowB += srcDescPtr->strides.hStride; + dstPtrRowR += dstDescPtr->strides.hStride; + dstPtrRowG += dstDescPtr->strides.hStride; + dstPtrRowB += dstDescPtr->strides.hStride; + } + } + } + + return RPP_SUCCESS; +} + +RppStatus color_temperature_f32_f32_host_tensor(Rpp32f *srcPtr, + RpptDescPtr srcDescPtr, + Rpp32f *dstPtr, + RpptDescPtr dstDescPtr, + Rpp8s *adjustmentValueTensor, + RpptROIPtr roiTensorPtrSrc, + RpptRoiType roiType, + RppLayoutParams layoutParams) +{ + RpptROI roiDefault = {0, 0, (Rpp32s)srcDescPtr->w, (Rpp32s)srcDescPtr->h}; + + omp_set_dynamic(0); +#pragma omp parallel for num_threads(dstDescPtr->n) + for(int batchCount = 0; batchCount < dstDescPtr->n; batchCount++) + { + RpptROI roi; + RpptROIPtr roiPtrInput = &roiTensorPtrSrc[batchCount]; + compute_roi_validation_host(roiPtrInput, &roi, &roiDefault, roiType); + + Rpp32f adjustmentValue = adjustmentValueTensor[batchCount] * ONE_OVER_255; + + Rpp32f *srcPtrImage, *dstPtrImage; + srcPtrImage = srcPtr + batchCount * srcDescPtr->strides.nStride; + dstPtrImage = dstPtr + batchCount * dstDescPtr->strides.nStride; + + Rpp32u bufferLength = roi.xywhROI.roiWidth * layoutParams.bufferMultiplier; + + __m256 pAdj = _mm256_set1_ps(adjustmentValue); + + Rpp32f *srcPtrChannel, *dstPtrChannel; + srcPtrChannel = srcPtrImage + (roi.xywhROI.xy.y * srcDescPtr->strides.hStride) + (roi.xywhROI.xy.x * layoutParams.bufferMultiplier); + dstPtrChannel = dstPtrImage; + + // Color Temperature with fused output-layout toggle (NHWC -> NCHW) + if ((srcDescPtr->c == 3) && (srcDescPtr->layout == RpptLayout::NHWC) && (dstDescPtr->layout == RpptLayout::NCHW)) + { + Rpp32u alignedLength = (bufferLength / 24) * 24; + + Rpp32f *srcPtrRow, *dstPtrRowR, *dstPtrRowG, *dstPtrRowB; + srcPtrRow = srcPtrChannel; + dstPtrRowR = dstPtrChannel; + dstPtrRowG = dstPtrRowR + dstDescPtr->strides.cStride; + dstPtrRowB = dstPtrRowG + dstDescPtr->strides.cStride; + + for(int i = 0; i < roi.xywhROI.roiHeight; i++) + { + Rpp32f *srcPtrTemp, *dstPtrTempR, *dstPtrTempG, *dstPtrTempB; + srcPtrTemp = srcPtrRow; + dstPtrTempR = dstPtrRowR; + dstPtrTempG = dstPtrRowG; + dstPtrTempB = dstPtrRowB; + + int vectorLoopCount = 0; + for (; vectorLoopCount < alignedLength; vectorLoopCount += 24) + { + __m256 p[3]; + + rpp_simd_load(rpp_load24_f32pkd3_to_f32pln3_avx, srcPtrTemp, p); // simd loads + compute_color_temperature_24_host(p, pAdj); // color_temperature adjustment + rpp_simd_store(rpp_store24_f32pln3_to_f32pln3_avx, dstPtrTempR, dstPtrTempG, dstPtrTempB, p); // simd stores + + srcPtrTemp += 24; + dstPtrTempR += 8; + dstPtrTempG += 8; + dstPtrTempB += 8; + } + for (; vectorLoopCount < bufferLength; vectorLoopCount += 3) + { + *dstPtrTempR++ = RPPPIXELCHECKF32(srcPtrTemp[0] + adjustmentValue); + *dstPtrTempG++ = RPPPIXELCHECKF32(srcPtrTemp[1]); + *dstPtrTempB++ = RPPPIXELCHECKF32(srcPtrTemp[2] - adjustmentValue); + + srcPtrTemp += 3; + } + + srcPtrRow += srcDescPtr->strides.hStride; + dstPtrRowR += dstDescPtr->strides.hStride; + dstPtrRowG += dstDescPtr->strides.hStride; + dstPtrRowB += dstDescPtr->strides.hStride; + } + } + + // Color Temperature with fused output-layout toggle (NCHW -> NHWC) + else if ((srcDescPtr->c == 3) && (srcDescPtr->layout == RpptLayout::NCHW) && (dstDescPtr->layout == RpptLayout::NHWC)) + { + Rpp32u alignedLength = (bufferLength / 24) * 24; + + Rpp32f *srcPtrRowR, *srcPtrRowG, *srcPtrRowB, *dstPtrRow; + srcPtrRowR = srcPtrChannel; + srcPtrRowG = srcPtrRowR + srcDescPtr->strides.cStride; + srcPtrRowB = srcPtrRowG + srcDescPtr->strides.cStride; + dstPtrRow = dstPtrChannel; + + for(int i = 0; i < roi.xywhROI.roiHeight; i++) + { + Rpp32f *srcPtrTempR, *srcPtrTempG, *srcPtrTempB, *dstPtrTemp; + srcPtrTempR = srcPtrRowR; + srcPtrTempG = srcPtrRowG; + srcPtrTempB = srcPtrRowB; + dstPtrTemp = dstPtrRow; + + int vectorLoopCount = 0; + for (; vectorLoopCount < alignedLength; vectorLoopCount += 8) + { + __m256 p[3]; + + rpp_simd_load(rpp_load24_f32pln3_to_f32pln3_avx, srcPtrTempR, srcPtrTempG, srcPtrTempB, p); // simd loads + compute_color_temperature_24_host(p, pAdj); // color_temperature adjustment + rpp_simd_store(rpp_store24_f32pln3_to_f32pkd3_avx, dstPtrTemp, p); // simd stores + + srcPtrTempR += 8; + srcPtrTempG += 8; + srcPtrTempB += 8; + dstPtrTemp += 24; + } + for (; vectorLoopCount < bufferLength; vectorLoopCount++) + { + dstPtrTemp[0] = RPPPIXELCHECKF32(*srcPtrTempR + adjustmentValue); + dstPtrTemp[1] = RPPPIXELCHECKF32(*srcPtrTempG); + dstPtrTemp[2] = RPPPIXELCHECKF32(*srcPtrTempB - adjustmentValue); + + dstPtrTemp += 3; + srcPtrTempR++; + srcPtrTempG++; + srcPtrTempB++; + } + + srcPtrRowR += srcDescPtr->strides.hStride; + srcPtrRowG += srcDescPtr->strides.hStride; + srcPtrRowB += srcDescPtr->strides.hStride; + dstPtrRow += dstDescPtr->strides.hStride; + } + } + + // Color Temperature with fused output-layout toggle (NHWC -> NHWC) + else if ((srcDescPtr->c == 3) && (srcDescPtr->layout == RpptLayout::NHWC) && (dstDescPtr->layout == RpptLayout::NHWC)) + { + Rpp32u alignedLength = (bufferLength / 24) * 24; + + Rpp32f *srcPtrRow, *dstPtrRow; + srcPtrRow = srcPtrChannel; + dstPtrRow = dstPtrChannel; + + for(int i = 0; i < roi.xywhROI.roiHeight; i++) + { + Rpp32f *srcPtrTemp, *dstPtrTemp; + srcPtrTemp = srcPtrRow; + dstPtrTemp = dstPtrRow; + + int vectorLoopCount = 0; + for (; vectorLoopCount < alignedLength; vectorLoopCount += 24) + { + __m256 p[3]; + + rpp_simd_load(rpp_load24_f32pkd3_to_f32pln3_avx, srcPtrTemp, p); // simd loads + compute_color_temperature_24_host(p, pAdj); // color_temperature adjustment + rpp_simd_store(rpp_store24_f32pln3_to_f32pkd3_avx, dstPtrTemp, p); // simd stores + + srcPtrTemp += 24; + dstPtrTemp += 24; + } + for (; vectorLoopCount < bufferLength; vectorLoopCount += 3) + { + dstPtrTemp[0] = RPPPIXELCHECKF32(srcPtrTemp[0] + adjustmentValue); + dstPtrTemp[1] = RPPPIXELCHECKF32(srcPtrTemp[1]); + dstPtrTemp[2] = RPPPIXELCHECKF32(srcPtrTemp[2] - adjustmentValue); + + srcPtrTemp += 3; + dstPtrTemp += 3; + } + + srcPtrRow += srcDescPtr->strides.hStride; + dstPtrRow += dstDescPtr->strides.hStride; + } + } + + // Color Temperature with fused output-layout toggle (NCHW -> NCHW) + else if ((srcDescPtr->c == 3) && (srcDescPtr->layout == RpptLayout::NCHW) && (dstDescPtr->layout == RpptLayout::NCHW)) + { + Rpp32u alignedLength = (bufferLength / 24) * 24; + + Rpp32f *srcPtrRowR, *srcPtrRowG, *srcPtrRowB, *dstPtrRowR, *dstPtrRowG, *dstPtrRowB; + srcPtrRowR = srcPtrChannel; + srcPtrRowG = srcPtrRowR + srcDescPtr->strides.cStride; + srcPtrRowB = srcPtrRowG + srcDescPtr->strides.cStride; + dstPtrRowR = dstPtrChannel; + dstPtrRowG = dstPtrRowR + dstDescPtr->strides.cStride; + dstPtrRowB = dstPtrRowG + dstDescPtr->strides.cStride; + + for(int i = 0; i < roi.xywhROI.roiHeight; i++) + { + Rpp32f *srcPtrTempR, *srcPtrTempG, *srcPtrTempB, *dstPtrTempR, *dstPtrTempG, *dstPtrTempB; + srcPtrTempR = srcPtrRowR; + srcPtrTempG = srcPtrRowG; + srcPtrTempB = srcPtrRowB; + dstPtrTempR = dstPtrRowR; + dstPtrTempG = dstPtrRowG; + dstPtrTempB = dstPtrRowB; + + int vectorLoopCount = 0; + for (; vectorLoopCount < alignedLength; vectorLoopCount += 8) + { + __m256 p[3]; + + rpp_simd_load(rpp_load24_f32pln3_to_f32pln3_avx, srcPtrTempR, srcPtrTempG, srcPtrTempB, p); // simd loads + compute_color_temperature_24_host(p, pAdj); // color_temperature adjustment + rpp_simd_store(rpp_store24_f32pln3_to_f32pln3_avx, dstPtrTempR, dstPtrTempG, dstPtrTempB, p); // simd stores + + srcPtrTempR += 8; + srcPtrTempG += 8; + srcPtrTempB += 8; + dstPtrTempR += 8; + dstPtrTempG += 8; + dstPtrTempB += 8; + } + for (; vectorLoopCount < bufferLength; vectorLoopCount++) + { + *dstPtrTempR++ = RPPPIXELCHECKF32(*srcPtrTempR + adjustmentValue); + *dstPtrTempG++ = RPPPIXELCHECKF32(*srcPtrTempG); + *dstPtrTempB++ = RPPPIXELCHECKF32(*srcPtrTempB - adjustmentValue); + + srcPtrTempR++; + srcPtrTempG++; + srcPtrTempB++; + } + + srcPtrRowR += srcDescPtr->strides.hStride; + srcPtrRowG += srcDescPtr->strides.hStride; + srcPtrRowB += srcDescPtr->strides.hStride; + dstPtrRowR += srcDescPtr->strides.hStride; + dstPtrRowG += srcDescPtr->strides.hStride; + dstPtrRowB += srcDescPtr->strides.hStride; + } + } + } + + return RPP_SUCCESS; +} + +RppStatus color_temperature_f16_f16_host_tensor(Rpp16f *srcPtr, + RpptDescPtr srcDescPtr, + Rpp16f *dstPtr, + RpptDescPtr dstDescPtr, + Rpp8s *adjustmentValueTensor, + RpptROIPtr roiTensorPtrSrc, + RpptRoiType roiType, + RppLayoutParams layoutParams) +{ + RpptROI roiDefault = {0, 0, (Rpp32s)srcDescPtr->w, (Rpp32s)srcDescPtr->h}; + + omp_set_dynamic(0); +#pragma omp parallel for num_threads(dstDescPtr->n) + for(int batchCount = 0; batchCount < dstDescPtr->n; batchCount++) + { + RpptROI roi; + RpptROIPtr roiPtrInput = &roiTensorPtrSrc[batchCount]; + compute_roi_validation_host(roiPtrInput, &roi, &roiDefault, roiType); + + Rpp32f adjustmentValue = adjustmentValueTensor[batchCount] * ONE_OVER_255; + + Rpp16f *srcPtrImage, *dstPtrImage; + srcPtrImage = srcPtr + batchCount * srcDescPtr->strides.nStride; + dstPtrImage = dstPtr + batchCount * dstDescPtr->strides.nStride; + + Rpp32u bufferLength = roi.xywhROI.roiWidth * layoutParams.bufferMultiplier; + + __m256 pAdj = _mm256_set1_ps(adjustmentValue); + + Rpp16f *srcPtrChannel, *dstPtrChannel; + srcPtrChannel = srcPtrImage + (roi.xywhROI.xy.y * srcDescPtr->strides.hStride) + (roi.xywhROI.xy.x * layoutParams.bufferMultiplier); + dstPtrChannel = dstPtrImage; + Rpp32u vectorIncrement = 24; + + // Color Temperature with fused output-layout toggle (NHWC -> NCHW) + if ((srcDescPtr->c == 3) && (srcDescPtr->layout == RpptLayout::NHWC) && (dstDescPtr->layout == RpptLayout::NCHW)) + { + Rpp32u alignedLength = (bufferLength / vectorIncrement) * vectorIncrement; + + Rpp16f *srcPtrRow, *dstPtrRowR, *dstPtrRowG, *dstPtrRowB; + srcPtrRow = srcPtrChannel; + dstPtrRowR = dstPtrChannel; + dstPtrRowG = dstPtrRowR + dstDescPtr->strides.cStride; + dstPtrRowB = dstPtrRowG + dstDescPtr->strides.cStride; + + for(int i = 0; i < roi.xywhROI.roiHeight; i++) + { + Rpp16f *srcPtrTemp, *dstPtrTempR, *dstPtrTempG, *dstPtrTempB; + srcPtrTemp = srcPtrRow; + dstPtrTempR = dstPtrRowR; + dstPtrTempG = dstPtrRowG; + dstPtrTempB = dstPtrRowB; + + int vectorLoopCount = 0; + for (; vectorLoopCount < alignedLength; vectorLoopCount += vectorIncrement) + { + Rpp32f srcPtrTemp_ps[24]; + Rpp32f dstPtrTempR_ps[8], dstPtrTempG_ps[8], dstPtrTempB_ps[8]; + + for(int cnt = 0; cnt < vectorIncrement; cnt++) + srcPtrTemp_ps[cnt] = (Rpp32f) srcPtrTemp[cnt]; + + __m256 p[3]; + + rpp_simd_load(rpp_load24_f32pkd3_to_f32pln3_avx, srcPtrTemp_ps, p); // simd loads + compute_color_temperature_24_host(p, pAdj); // color_temperature adjustment + rpp_simd_store(rpp_store24_f32pln3_to_f32pln3_avx, dstPtrTempR_ps, dstPtrTempG_ps, dstPtrTempB_ps, p); // simd stores + + for(int cnt = 0; cnt < 8; cnt++) + { + dstPtrTempR[cnt] = (Rpp16f) dstPtrTempR_ps[cnt]; + dstPtrTempG[cnt] = (Rpp16f) dstPtrTempG_ps[cnt]; + dstPtrTempB[cnt] = (Rpp16f) dstPtrTempB_ps[cnt]; + } + + srcPtrTemp += 24; + dstPtrTempR += 8; + dstPtrTempG += 8; + dstPtrTempB += 8; + } + for (; vectorLoopCount < bufferLength; vectorLoopCount += 3) + { + *dstPtrTempR++ = (Rpp16f) RPPPIXELCHECKF32(srcPtrTemp[0] + adjustmentValue); + *dstPtrTempG++ = (Rpp16f) RPPPIXELCHECKF32(srcPtrTemp[1]); + *dstPtrTempB++ = (Rpp16f) RPPPIXELCHECKF32(srcPtrTemp[2] - adjustmentValue); + + srcPtrTemp += 3; + } + + srcPtrRow += srcDescPtr->strides.hStride; + dstPtrRowR += dstDescPtr->strides.hStride; + dstPtrRowG += dstDescPtr->strides.hStride; + dstPtrRowB += dstDescPtr->strides.hStride; + } + } + + // Color Temperature with fused output-layout toggle (NCHW -> NHWC) + else if ((srcDescPtr->c == 3) && (srcDescPtr->layout == RpptLayout::NCHW) && (dstDescPtr->layout == RpptLayout::NHWC)) + { + Rpp32u alignedLength = (bufferLength / 24) * 24; + + Rpp16f *srcPtrRowR, *srcPtrRowG, *srcPtrRowB, *dstPtrRow; + srcPtrRowR = srcPtrChannel; + srcPtrRowG = srcPtrRowR + srcDescPtr->strides.cStride; + srcPtrRowB = srcPtrRowG + srcDescPtr->strides.cStride; + dstPtrRow = dstPtrChannel; + + for(int i = 0; i < roi.xywhROI.roiHeight; i++) + { + Rpp16f *srcPtrTempR, *srcPtrTempG, *srcPtrTempB, *dstPtrTemp; + srcPtrTempR = srcPtrRowR; + srcPtrTempG = srcPtrRowG; + srcPtrTempB = srcPtrRowB; + dstPtrTemp = dstPtrRow; + + int vectorLoopCount = 0; + for (; vectorLoopCount < alignedLength; vectorLoopCount += 8) + { + Rpp32f srcPtrTempR_ps[8], srcPtrTempG_ps[8], srcPtrTempB_ps[8]; + Rpp32f dstPtrTemp_ps[25]; + + for(int cnt = 0; cnt < 8; cnt++) + { + srcPtrTempR_ps[cnt] = (Rpp32f) srcPtrTempR[cnt]; + srcPtrTempG_ps[cnt] = (Rpp32f) srcPtrTempG[cnt]; + srcPtrTempB_ps[cnt] = (Rpp32f) srcPtrTempB[cnt]; + } + + __m256 p[3]; + + rpp_simd_load(rpp_load24_f32pln3_to_f32pln3_avx, srcPtrTempR_ps, srcPtrTempG_ps, srcPtrTempB_ps, p); // simd loads + compute_color_temperature_24_host(p, pAdj); // color_temperature adjustment + rpp_simd_store(rpp_store24_f32pln3_to_f32pkd3_avx, dstPtrTemp_ps, p); // simd stores + + for(int cnt = 0; cnt < 24; cnt++) + dstPtrTemp[cnt] = (Rpp16f) dstPtrTemp_ps[cnt]; + + srcPtrTempR += 8; + srcPtrTempG += 8; + srcPtrTempB += 8; + dstPtrTemp += 24; + } + for (; vectorLoopCount < bufferLength; vectorLoopCount++) + { + dstPtrTemp[0] = (Rpp16f) RPPPIXELCHECKF32(*srcPtrTempR + adjustmentValue); + dstPtrTemp[1] = (Rpp16f) RPPPIXELCHECKF32(*srcPtrTempG); + dstPtrTemp[2] = (Rpp16f) RPPPIXELCHECKF32(*srcPtrTempB - adjustmentValue); + + dstPtrTemp += 3; + srcPtrTempR++; + srcPtrTempG++; + srcPtrTempB++; + } + + srcPtrRowR += srcDescPtr->strides.hStride; + srcPtrRowG += srcDescPtr->strides.hStride; + srcPtrRowB += srcDescPtr->strides.hStride; + dstPtrRow += dstDescPtr->strides.hStride; + } + } + + // Color Temperature with fused output-layout toggle (NHWC -> NHWC) + else if ((srcDescPtr->c == 3) && (srcDescPtr->layout == RpptLayout::NHWC) && (dstDescPtr->layout == RpptLayout::NHWC)) + { + Rpp32u alignedLength = (bufferLength / 24) * 24; + + Rpp16f *srcPtrRow, *dstPtrRow; + srcPtrRow = srcPtrChannel; + dstPtrRow = dstPtrChannel; + + for(int i = 0; i < roi.xywhROI.roiHeight; i++) + { + Rpp16f *srcPtrTemp, *dstPtrTemp; + srcPtrTemp = srcPtrRow; + dstPtrTemp = dstPtrRow; + + int vectorLoopCount = 0; + for (; vectorLoopCount < alignedLength; vectorLoopCount += 24) + { + Rpp32f srcPtrTemp_ps[24], dstPtrTemp_ps[25]; + + for(int cnt = 0; cnt < 24; cnt++) + srcPtrTemp_ps[cnt] = (Rpp32f) srcPtrTemp[cnt]; + + __m256 p[3]; + + rpp_simd_load(rpp_load24_f32pkd3_to_f32pln3_avx, srcPtrTemp_ps, p); // simd loads + compute_color_temperature_24_host(p, pAdj); // color_temperature adjustment + rpp_simd_store(rpp_store24_f32pln3_to_f32pkd3_avx, dstPtrTemp_ps, p); // simd stores + + for(int cnt = 0; cnt < 24; cnt++) + dstPtrTemp[cnt] = (Rpp16f) dstPtrTemp_ps[cnt]; + + srcPtrTemp += 24; + dstPtrTemp += 24; + } + for (; vectorLoopCount < bufferLength; vectorLoopCount += 3) + { + dstPtrTemp[0] = (Rpp16f) RPPPIXELCHECKF32(srcPtrTemp[0] + adjustmentValue); + dstPtrTemp[1] = (Rpp16f) RPPPIXELCHECKF32(srcPtrTemp[1]); + dstPtrTemp[2] = (Rpp16f) RPPPIXELCHECKF32(srcPtrTemp[2] - adjustmentValue); + + srcPtrTemp += 3; + dstPtrTemp += 3; + } + + srcPtrRow += srcDescPtr->strides.hStride; + dstPtrRow += dstDescPtr->strides.hStride; + } + } + + // Color Temperature with fused output-layout toggle (NCHW -> NCHW) + else if ((srcDescPtr->c == 3) && (srcDescPtr->layout == RpptLayout::NCHW) && (dstDescPtr->layout == RpptLayout::NCHW)) + { + Rpp32u alignedLength = (bufferLength / 24) * 24; + + Rpp16f *srcPtrRowR, *srcPtrRowG, *srcPtrRowB, *dstPtrRowR, *dstPtrRowG, *dstPtrRowB; + srcPtrRowR = srcPtrChannel; + srcPtrRowG = srcPtrRowR + srcDescPtr->strides.cStride; + srcPtrRowB = srcPtrRowG + srcDescPtr->strides.cStride; + dstPtrRowR = dstPtrChannel; + dstPtrRowG = dstPtrRowR + dstDescPtr->strides.cStride; + dstPtrRowB = dstPtrRowG + dstDescPtr->strides.cStride; + + for(int i = 0; i < roi.xywhROI.roiHeight; i++) + { + Rpp16f *srcPtrTempR, *srcPtrTempG, *srcPtrTempB, *dstPtrTempR, *dstPtrTempG, *dstPtrTempB; + srcPtrTempR = srcPtrRowR; + srcPtrTempG = srcPtrRowG; + srcPtrTempB = srcPtrRowB; + dstPtrTempR = dstPtrRowR; + dstPtrTempG = dstPtrRowG; + dstPtrTempB = dstPtrRowB; + + int vectorLoopCount = 0; + for (; vectorLoopCount < alignedLength; vectorLoopCount += 8) + { + Rpp32f srcPtrTempR_ps[8], srcPtrTempG_ps[8], srcPtrTempB_ps[8]; + Rpp32f dstPtrTempR_ps[8], dstPtrTempG_ps[8], dstPtrTempB_ps[8]; + + for(int cnt = 0; cnt < 8; cnt++) + { + srcPtrTempR_ps[cnt] = (Rpp32f) srcPtrTempR[cnt]; + srcPtrTempG_ps[cnt] = (Rpp32f) srcPtrTempG[cnt]; + srcPtrTempB_ps[cnt] = (Rpp32f) srcPtrTempB[cnt]; + } + + __m256 p[3]; + + rpp_simd_load(rpp_load24_f32pln3_to_f32pln3_avx, srcPtrTempR_ps, srcPtrTempG_ps, srcPtrTempB_ps, p); // simd loads + compute_color_temperature_24_host(p, pAdj); // color_temperature adjustment + rpp_simd_store(rpp_store24_f32pln3_to_f32pln3_avx, dstPtrTempR_ps, dstPtrTempG_ps, dstPtrTempB_ps, p); // simd stores + + for(int cnt = 0; cnt < 8; cnt++) + { + dstPtrTempR[cnt] = (Rpp16f) dstPtrTempR_ps[cnt]; + dstPtrTempG[cnt] = (Rpp16f) dstPtrTempG_ps[cnt]; + dstPtrTempB[cnt] = (Rpp16f) dstPtrTempB_ps[cnt]; + } + + srcPtrTempR += 8; + srcPtrTempG += 8; + srcPtrTempB += 8; + dstPtrTempR += 8; + dstPtrTempG += 8; + dstPtrTempB += 8; + } + for (; vectorLoopCount < bufferLength; vectorLoopCount++) + { + *dstPtrTempR++ = (Rpp16f) RPPPIXELCHECKF32(*srcPtrTempR + adjustmentValue); + *dstPtrTempG++ = (Rpp16f) RPPPIXELCHECKF32(*srcPtrTempG); + *dstPtrTempB++ = (Rpp16f) RPPPIXELCHECKF32(*srcPtrTempB - adjustmentValue); + + srcPtrTempR++; + srcPtrTempG++; + srcPtrTempB++; + } + + srcPtrRowR += srcDescPtr->strides.hStride; + srcPtrRowG += srcDescPtr->strides.hStride; + srcPtrRowB += srcDescPtr->strides.hStride; + dstPtrRowR += srcDescPtr->strides.hStride; + dstPtrRowG += srcDescPtr->strides.hStride; + dstPtrRowB += srcDescPtr->strides.hStride; + } + } + } + + return RPP_SUCCESS; +} + +RppStatus color_temperature_i8_i8_host_tensor(Rpp8s *srcPtr, + RpptDescPtr srcDescPtr, + Rpp8s *dstPtr, + RpptDescPtr dstDescPtr, + Rpp8s *adjustmentValueTensor, + RpptROIPtr roiTensorPtrSrc, + RpptRoiType roiType, + RppLayoutParams layoutParams) +{ + RpptROI roiDefault = {0, 0, (Rpp32s)srcDescPtr->w, (Rpp32s)srcDescPtr->h}; + + omp_set_dynamic(0); +#pragma omp parallel for num_threads(dstDescPtr->n) + for(int batchCount = 0; batchCount < dstDescPtr->n; batchCount++) + { + RpptROI roi; + RpptROIPtr roiPtrInput = &roiTensorPtrSrc[batchCount]; + compute_roi_validation_host(roiPtrInput, &roi, &roiDefault, roiType); + + Rpp32f adjustmentValue = adjustmentValueTensor[batchCount]; + + Rpp8s *srcPtrImage, *dstPtrImage; + srcPtrImage = srcPtr + batchCount * srcDescPtr->strides.nStride; + dstPtrImage = dstPtr + batchCount * dstDescPtr->strides.nStride; + + Rpp32u bufferLength = roi.xywhROI.roiWidth * layoutParams.bufferMultiplier; + + __m256 pAdj = _mm256_set1_ps(adjustmentValue); + + Rpp8s *srcPtrChannel, *dstPtrChannel; + srcPtrChannel = srcPtrImage + (roi.xywhROI.xy.y * srcDescPtr->strides.hStride) + (roi.xywhROI.xy.x * layoutParams.bufferMultiplier); + dstPtrChannel = dstPtrImage; + + // Color Temperature with fused output-layout toggle (NHWC -> NCHW) + if ((srcDescPtr->c == 3) && (srcDescPtr->layout == RpptLayout::NHWC) && (dstDescPtr->layout == RpptLayout::NCHW)) + { + Rpp32u alignedLength = (bufferLength / 48) * 48; + + Rpp8s *srcPtrRow, *dstPtrRowR, *dstPtrRowG, *dstPtrRowB; + srcPtrRow = srcPtrChannel; + dstPtrRowR = dstPtrChannel; + dstPtrRowG = dstPtrRowR + dstDescPtr->strides.cStride; + dstPtrRowB = dstPtrRowG + dstDescPtr->strides.cStride; + + for(int i = 0; i < roi.xywhROI.roiHeight; i++) + { + Rpp8s *srcPtrTemp, *dstPtrTempR, *dstPtrTempG, *dstPtrTempB; + srcPtrTemp = srcPtrRow; + dstPtrTempR = dstPtrRowR; + dstPtrTempG = dstPtrRowG; + dstPtrTempB = dstPtrRowB; + + int vectorLoopCount = 0; + for (; vectorLoopCount < alignedLength; vectorLoopCount += 48) + { + __m256 p[6]; + + rpp_simd_load(rpp_load48_i8pkd3_to_f32pln3_avx, srcPtrTemp, p); // simd loads + compute_color_temperature_48_host(p, pAdj); // color_temperature adjustment + rpp_simd_store(rpp_store48_f32pln3_to_i8pln3_avx, dstPtrTempR, dstPtrTempG, dstPtrTempB, p); // simd stores + + srcPtrTemp += 48; + dstPtrTempR += 16; + dstPtrTempG += 16; + dstPtrTempB += 16; + } + for (; vectorLoopCount < bufferLength; vectorLoopCount += 3) + { + *dstPtrTempR++ = (Rpp8s) RPPPIXELCHECKI8(srcPtrTemp[0] + adjustmentValue); + *dstPtrTempG++ = (Rpp8s) RPPPIXELCHECKI8(srcPtrTemp[1]); + *dstPtrTempB++ = (Rpp8s) RPPPIXELCHECKI8(srcPtrTemp[2] - adjustmentValue); + + srcPtrTemp += 3; + } + + srcPtrRow += srcDescPtr->strides.hStride; + dstPtrRowR += dstDescPtr->strides.hStride; + dstPtrRowG += dstDescPtr->strides.hStride; + dstPtrRowB += dstDescPtr->strides.hStride; + } + } + + // Color Temperature with fused output-layout toggle (NCHW -> NHWC) + else if ((srcDescPtr->c == 3) && (srcDescPtr->layout == RpptLayout::NCHW) && (dstDescPtr->layout == RpptLayout::NHWC)) + { + Rpp32u alignedLength = (bufferLength / 48) * 48; + + Rpp8s *srcPtrRowR, *srcPtrRowG, *srcPtrRowB, *dstPtrRow; + srcPtrRowR = srcPtrChannel; + srcPtrRowG = srcPtrRowR + srcDescPtr->strides.cStride; + srcPtrRowB = srcPtrRowG + srcDescPtr->strides.cStride; + dstPtrRow = dstPtrChannel; + + for(int i = 0; i < roi.xywhROI.roiHeight; i++) + { + Rpp8s *srcPtrTempR, *srcPtrTempG, *srcPtrTempB, *dstPtrTemp; + srcPtrTempR = srcPtrRowR; + srcPtrTempG = srcPtrRowG; + srcPtrTempB = srcPtrRowB; + dstPtrTemp = dstPtrRow; + + int vectorLoopCount = 0; + for (; vectorLoopCount < alignedLength; vectorLoopCount += 16) + { + __m256 p[6]; + + rpp_simd_load(rpp_load48_i8pln3_to_f32pln3_avx, srcPtrTempR, srcPtrTempG, srcPtrTempB, p); // simd loads + compute_color_temperature_48_host(p, pAdj); // color_temperature adjustment + rpp_simd_store(rpp_store48_f32pln3_to_i8pkd3_avx, dstPtrTemp, p); // simd stores + + srcPtrTempR += 16; + srcPtrTempG += 16; + srcPtrTempB += 16; + dstPtrTemp += 48; + } + for (; vectorLoopCount < bufferLength; vectorLoopCount++) + { + dstPtrTemp[0] = (Rpp8s) RPPPIXELCHECKI8(*srcPtrTempR + adjustmentValue); + dstPtrTemp[1] = (Rpp8s) RPPPIXELCHECKI8(*srcPtrTempG); + dstPtrTemp[2] = (Rpp8s) RPPPIXELCHECKI8(*srcPtrTempB - adjustmentValue); + + dstPtrTemp += 3; + srcPtrTempR++; + srcPtrTempG++; + srcPtrTempB++; + } + + srcPtrRowR += srcDescPtr->strides.hStride; + srcPtrRowG += srcDescPtr->strides.hStride; + srcPtrRowB += srcDescPtr->strides.hStride; + dstPtrRow += dstDescPtr->strides.hStride; + } + } + + // Color Temperature with fused output-layout toggle (NHWC -> NHWC) + else if ((srcDescPtr->c == 3) && (srcDescPtr->layout == RpptLayout::NHWC) && (dstDescPtr->layout == RpptLayout::NHWC)) + { + Rpp32u alignedLength = (bufferLength / 48) * 48; + + Rpp8s *srcPtrRow, *dstPtrRow; + srcPtrRow = srcPtrChannel; + dstPtrRow = dstPtrChannel; + + for(int i = 0; i < roi.xywhROI.roiHeight; i++) + { + Rpp8s *srcPtrTemp, *dstPtrTemp; + srcPtrTemp = srcPtrRow; + dstPtrTemp = dstPtrRow; + + int vectorLoopCount = 0; + for (; vectorLoopCount < alignedLength; vectorLoopCount += 48) + { + __m256 p[6]; + + rpp_simd_load(rpp_load48_i8pkd3_to_f32pln3_avx, srcPtrTemp, p); // simd loads + compute_color_temperature_48_host(p, pAdj); // color_temperature adjustment + rpp_simd_store(rpp_store48_f32pln3_to_i8pkd3_avx, dstPtrTemp, p); // simd stores + + srcPtrTemp += 48; + dstPtrTemp += 48; + } + for (; vectorLoopCount < bufferLength; vectorLoopCount += 3) + { + dstPtrTemp[0] = (Rpp8s) RPPPIXELCHECKI8(srcPtrTemp[0] + adjustmentValue); + dstPtrTemp[1] = (Rpp8s) RPPPIXELCHECKI8(srcPtrTemp[1]); + dstPtrTemp[2] = (Rpp8s) RPPPIXELCHECKI8(srcPtrTemp[2] - adjustmentValue); + + srcPtrTemp += 3; + dstPtrTemp += 3; + } + + srcPtrRow += srcDescPtr->strides.hStride; + dstPtrRow += dstDescPtr->strides.hStride; + } + } + + // Color Temperature with fused output-layout toggle (NCHW -> NCHW) + else if ((srcDescPtr->c == 3) && (srcDescPtr->layout == RpptLayout::NCHW) && (dstDescPtr->layout == RpptLayout::NCHW)) + { + Rpp32u alignedLength = (bufferLength / 48) * 48; + + Rpp8s *srcPtrRowR, *srcPtrRowG, *srcPtrRowB, *dstPtrRowR, *dstPtrRowG, *dstPtrRowB; + srcPtrRowR = srcPtrChannel; + srcPtrRowG = srcPtrRowR + srcDescPtr->strides.cStride; + srcPtrRowB = srcPtrRowG + srcDescPtr->strides.cStride; + dstPtrRowR = dstPtrChannel; + dstPtrRowG = dstPtrRowR + dstDescPtr->strides.cStride; + dstPtrRowB = dstPtrRowG + dstDescPtr->strides.cStride; + + for(int i = 0; i < roi.xywhROI.roiHeight; i++) + { + Rpp8s *srcPtrTempR, *srcPtrTempG, *srcPtrTempB, *dstPtrTempR, *dstPtrTempG, *dstPtrTempB; + srcPtrTempR = srcPtrRowR; + srcPtrTempG = srcPtrRowG; + srcPtrTempB = srcPtrRowB; + dstPtrTempR = dstPtrRowR; + dstPtrTempG = dstPtrRowG; + dstPtrTempB = dstPtrRowB; + + int vectorLoopCount = 0; + for (; vectorLoopCount < alignedLength; vectorLoopCount += 16) + { + __m256 p[6]; + + rpp_simd_load(rpp_load48_i8pln3_to_f32pln3_avx, srcPtrTempR, srcPtrTempG, srcPtrTempB, p); // simd loads + compute_color_temperature_48_host(p, pAdj); // color_temperature adjustment + rpp_simd_store(rpp_store48_f32pln3_to_i8pln3_avx, dstPtrTempR, dstPtrTempG, dstPtrTempB, p); // simd stores + + srcPtrTempR += 16; + srcPtrTempG += 16; + srcPtrTempB += 16; + dstPtrTempR += 16; + dstPtrTempG += 16; + dstPtrTempB += 16; + } + for (; vectorLoopCount < bufferLength; vectorLoopCount++) + { + *dstPtrTempR++ = (Rpp8s) RPPPIXELCHECKI8(*srcPtrTempR + adjustmentValue); + *dstPtrTempG++ = (Rpp8s) RPPPIXELCHECKI8(*srcPtrTempG); + *dstPtrTempB++ = (Rpp8s) RPPPIXELCHECKI8(*srcPtrTempB - adjustmentValue); + + srcPtrTempR++; + srcPtrTempG++; + srcPtrTempB++; + } + + srcPtrRowR += srcDescPtr->strides.hStride; + srcPtrRowG += srcDescPtr->strides.hStride; + srcPtrRowB += srcDescPtr->strides.hStride; + dstPtrRowR += dstDescPtr->strides.hStride; + dstPtrRowG += dstDescPtr->strides.hStride; + dstPtrRowB += dstDescPtr->strides.hStride; + } + } + } + + return RPP_SUCCESS; +} diff --git a/src/modules/hip/hip_tensor_color_augmentations.hpp b/src/modules/hip/hip_tensor_color_augmentations.hpp index 873f06b97..c5610dbcb 100644 --- a/src/modules/hip/hip_tensor_color_augmentations.hpp +++ b/src/modules/hip/hip_tensor_color_augmentations.hpp @@ -33,5 +33,6 @@ SOFTWARE. #include "kernel/exposure.hpp" #include "kernel/contrast.hpp" #include "kernel/lut.hpp" +#include "kernel/color_temperature.hpp" #endif // HIP_TENSOR_COLOR_AUGMENTATIONS_HPP diff --git a/src/modules/hip/kernel/color_temperature.hpp b/src/modules/hip/kernel/color_temperature.hpp new file mode 100644 index 000000000..ad8adc32a --- /dev/null +++ b/src/modules/hip/kernel/color_temperature.hpp @@ -0,0 +1,223 @@ +#include +#include "rpp_hip_common.hpp" + +template +__device__ void color_temperature_hip_compute(T *srcPtr, d_float24 *pix_f24, float4 *adjustmentValue_f4) +{ + float4 adjustment_f4; + if constexpr ((std::is_same::value) || (std::is_same::value)) + { + adjustment_f4 = *adjustmentValue_f4 * (float4) ONE_OVER_255; + rpp_hip_math_add8_const(&pix_f24->f8[0], &pix_f24->f8[0], adjustment_f4); + rpp_hip_math_subtract8_const(&pix_f24->f8[2], &pix_f24->f8[2], adjustment_f4); + } + else if constexpr (std::is_same::value) + { + adjustment_f4 = *adjustmentValue_f4; + rpp_hip_math_add24_const(pix_f24, pix_f24, (float4)128); + rpp_hip_math_add8_const(&pix_f24->f8[0], &pix_f24->f8[0], adjustment_f4); + rpp_hip_math_subtract8_const(&pix_f24->f8[2], &pix_f24->f8[2], adjustment_f4); + rpp_hip_pixel_check_0to255(pix_f24); + rpp_hip_math_subtract24_const(pix_f24, pix_f24, (float4)128); + } + else + { + rpp_hip_math_add8_const(&pix_f24->f8[0], &pix_f24->f8[0], *adjustmentValue_f4); + rpp_hip_math_subtract8_const(&pix_f24->f8[2], &pix_f24->f8[2], *adjustmentValue_f4); + } +} + +template +__global__ void color_temperature_pkd_hip_tensor(T *srcPtr, + uint2 srcStridesNH, + T *dstPtr, + uint2 dstStridesNH, + int *adjustmentValueTensor, + RpptROIPtr roiTensorPtrSrc) +{ + int id_x = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x) * 8; + int id_y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y; + int id_z = hipBlockIdx_z * hipBlockDim_z + hipThreadIdx_z; + + if ((id_y >= roiTensorPtrSrc[id_z].xywhROI.roiHeight) || (id_x >= roiTensorPtrSrc[id_z].xywhROI.roiWidth)) + { + return; + } + + uint srcIdx = (id_z * srcStridesNH.x) + ((id_y + roiTensorPtrSrc[id_z].xywhROI.xy.y) * srcStridesNH.y) + ((id_x + roiTensorPtrSrc[id_z].xywhROI.xy.x) * 3); + uint dstIdx = (id_z * dstStridesNH.x) + (id_y * dstStridesNH.y) + id_x * 3; + + float4 adjustmentValue_f4 = (float4)((float)adjustmentValueTensor[id_z]); + + d_float24 pix_f24; + + rpp_hip_load24_pkd3_and_unpack_to_float24_pln3(srcPtr + srcIdx, &pix_f24); + color_temperature_hip_compute(srcPtr, &pix_f24, &adjustmentValue_f4); + rpp_hip_pack_float24_pln3_and_store24_pkd3(dstPtr + dstIdx, &pix_f24); +} + +template +__global__ void color_temperature_pln_hip_tensor(T *srcPtr, + uint3 srcStridesNCH, + T *dstPtr, + uint3 dstStridesNCH, + int *adjustmentValueTensor, + RpptROIPtr roiTensorPtrSrc) +{ + int id_x = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x) * 8; + int id_y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y; + int id_z = hipBlockIdx_z * hipBlockDim_z + hipThreadIdx_z; + + if ((id_y >= roiTensorPtrSrc[id_z].xywhROI.roiHeight) || (id_x >= roiTensorPtrSrc[id_z].xywhROI.roiWidth)) + { + return; + } + + uint srcIdx = (id_z * srcStridesNCH.x) + ((id_y + roiTensorPtrSrc[id_z].xywhROI.xy.y) * srcStridesNCH.z) + (id_x + roiTensorPtrSrc[id_z].xywhROI.xy.x); + uint dstIdx = (id_z * dstStridesNCH.x) + (id_y * dstStridesNCH.z) + id_x; + + float4 adjustmentValue_f4 = (float4)((float)adjustmentValueTensor[id_z]); + + d_float24 pix_f24; + + rpp_hip_load24_pln3_and_unpack_to_float24_pln3(srcPtr + srcIdx, srcStridesNCH.y, &pix_f24); + color_temperature_hip_compute(srcPtr, &pix_f24, &adjustmentValue_f4); + rpp_hip_pack_float24_pln3_and_store24_pln3(dstPtr + dstIdx, dstStridesNCH.y, &pix_f24); +} + +template +__global__ void color_temperature_pkd3_pln3_hip_tensor(T *srcPtr, + uint2 srcStridesNH, + T *dstPtr, + uint3 dstStridesNCH, + int *adjustmentValueTensor, + RpptROIPtr roiTensorPtrSrc) +{ + int id_x = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x) * 8; + int id_y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y; + int id_z = hipBlockIdx_z * hipBlockDim_z + hipThreadIdx_z; + + if ((id_y >= roiTensorPtrSrc[id_z].xywhROI.roiHeight) || (id_x >= roiTensorPtrSrc[id_z].xywhROI.roiWidth)) + { + return; + } + + uint srcIdx = (id_z * srcStridesNH.x) + ((id_y + roiTensorPtrSrc[id_z].xywhROI.xy.y) * srcStridesNH.y) + ((id_x + roiTensorPtrSrc[id_z].xywhROI.xy.x) * 3); + uint dstIdx = (id_z * dstStridesNCH.x) + (id_y * dstStridesNCH.z) + id_x; + + float4 adjustmentValue_f4 = (float4)((float)adjustmentValueTensor[id_z]); + + d_float24 pix_f24; + + rpp_hip_load24_pkd3_and_unpack_to_float24_pln3(srcPtr + srcIdx, &pix_f24); + color_temperature_hip_compute(srcPtr, &pix_f24, &adjustmentValue_f4); + rpp_hip_pack_float24_pln3_and_store24_pln3(dstPtr + dstIdx, dstStridesNCH.y, &pix_f24); +} + +template +__global__ void color_temperature_pln3_pkd3_hip_tensor(T *srcPtr, + uint3 srcStridesNCH, + T *dstPtr, + uint2 dstStridesNH, + int *adjustmentValueTensor, + RpptROIPtr roiTensorPtrSrc) +{ + int id_x = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x) * 8; + int id_y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y; + int id_z = hipBlockIdx_z * hipBlockDim_z + hipThreadIdx_z; + + if ((id_y >= roiTensorPtrSrc[id_z].xywhROI.roiHeight) || (id_x >= roiTensorPtrSrc[id_z].xywhROI.roiWidth)) + { + return; + } + + uint srcIdx = (id_z * srcStridesNCH.x) + ((id_y + roiTensorPtrSrc[id_z].xywhROI.xy.y) * srcStridesNCH.z) + (id_x + roiTensorPtrSrc[id_z].xywhROI.xy.x); + uint dstIdx = (id_z * dstStridesNH.x) + (id_y * dstStridesNH.y) + id_x * 3; + + float4 adjustmentValue_f4 = (float4)((float)adjustmentValueTensor[id_z]); + + d_float24 pix_f24; + + rpp_hip_load24_pln3_and_unpack_to_float24_pln3(srcPtr + srcIdx, srcStridesNCH.y, &pix_f24); + color_temperature_hip_compute(srcPtr, &pix_f24, &adjustmentValue_f4); + rpp_hip_pack_float24_pln3_and_store24_pkd3(dstPtr + dstIdx, &pix_f24); +} + +template +RppStatus hip_exec_color_temperature_tensor(T *srcPtr, + RpptDescPtr srcDescPtr, + T *dstPtr, + RpptDescPtr dstDescPtr, + RpptROIPtr roiTensorPtrSrc, + RpptRoiType roiType, + rpp::Handle& handle) +{ + if (roiType == RpptRoiType::LTRB) + hip_exec_roi_converison_ltrb_to_xywh(roiTensorPtrSrc, handle); + + if ((srcDescPtr->c == 3) && (dstDescPtr->c == 3)) + { + int globalThreads_x = (dstDescPtr->strides.hStride + 7) >> 3; + int globalThreads_y = dstDescPtr->h; + int globalThreads_z = handle.GetBatchSize(); + + if ((srcDescPtr->layout == RpptLayout::NHWC) && (dstDescPtr->layout == RpptLayout::NHWC)) + { + hipLaunchKernelGGL(color_temperature_pkd_hip_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), + handle.GetInitHandle()->mem.mgpu.intArr[0].intmem, + roiTensorPtrSrc); + } + else if ((srcDescPtr->layout == RpptLayout::NCHW) && (dstDescPtr->layout == RpptLayout::NCHW)) + { + hipLaunchKernelGGL(color_temperature_pln_hip_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_uint3(srcDescPtr->strides.nStride, srcDescPtr->strides.cStride, srcDescPtr->strides.hStride), + dstPtr, + make_uint3(dstDescPtr->strides.nStride, dstDescPtr->strides.cStride, dstDescPtr->strides.hStride), + handle.GetInitHandle()->mem.mgpu.intArr[0].intmem, + roiTensorPtrSrc); + } + else if ((srcDescPtr->layout == RpptLayout::NHWC) && (dstDescPtr->layout == RpptLayout::NCHW)) + { + hipLaunchKernelGGL(color_temperature_pkd3_pln3_hip_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_uint3(dstDescPtr->strides.nStride, dstDescPtr->strides.cStride, dstDescPtr->strides.hStride), + handle.GetInitHandle()->mem.mgpu.intArr[0].intmem, + roiTensorPtrSrc); + } + else if ((srcDescPtr->layout == RpptLayout::NCHW) && (dstDescPtr->layout == RpptLayout::NHWC)) + { + hipLaunchKernelGGL(color_temperature_pln3_pkd3_hip_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_uint3(srcDescPtr->strides.nStride, srcDescPtr->strides.cStride, srcDescPtr->strides.hStride), + dstPtr, + make_uint2(dstDescPtr->strides.nStride, dstDescPtr->strides.hStride), + handle.GetInitHandle()->mem.mgpu.intArr[0].intmem, + roiTensorPtrSrc); + } + } + + return RPP_SUCCESS; +} \ No newline at end of file diff --git a/src/modules/rppt_tensor_color_augmentations.cpp b/src/modules/rppt_tensor_color_augmentations.cpp index be61b6da1..3023973fc 100644 --- a/src/modules/rppt_tensor_color_augmentations.cpp +++ b/src/modules/rppt_tensor_color_augmentations.cpp @@ -411,7 +411,7 @@ RppStatus rppt_color_cast_host(RppPtr_t srcPtr, { if (srcDescPtr->c != 3) { - return RPP_ERROR_INVALID_ARGUMENTS; + return RPP_ERROR_INVALID_CHANNELS; } RppLayoutParams layoutParams = get_layout_params(srcDescPtr->layout, srcDescPtr->c); @@ -671,6 +671,72 @@ RppStatus rppt_lut_host(RppPtr_t srcPtr, return RPP_SUCCESS; } +/******************** color_temperature ********************/ + +RppStatus rppt_color_temperature_host(RppPtr_t srcPtr, + RpptDescPtr srcDescPtr, + RppPtr_t dstPtr, + RpptDescPtr dstDescPtr, + Rpp8s *adjustmentValueTensor, + RpptROIPtr roiTensorPtrSrc, + RpptRoiType roiType, + rppHandle_t rppHandle) +{ + if (srcDescPtr->c != 3) + { + return RPP_ERROR_INVALID_CHANNELS; + } + + RppLayoutParams layoutParams = get_layout_params(srcDescPtr->layout, srcDescPtr->c); + + if ((srcDescPtr->dataType == RpptDataType::U8) && (dstDescPtr->dataType == RpptDataType::U8)) + { + color_temperature_u8_u8_host_tensor(static_cast(srcPtr) + srcDescPtr->offsetInBytes, + srcDescPtr, + static_cast(dstPtr) + dstDescPtr->offsetInBytes, + dstDescPtr, + adjustmentValueTensor, + roiTensorPtrSrc, + roiType, + layoutParams); + } + else if ((srcDescPtr->dataType == RpptDataType::F16) && (dstDescPtr->dataType == RpptDataType::F16)) + { + color_temperature_f16_f16_host_tensor(reinterpret_cast(static_cast(srcPtr) + srcDescPtr->offsetInBytes), + srcDescPtr, + reinterpret_cast(static_cast(dstPtr) + dstDescPtr->offsetInBytes), + dstDescPtr, + adjustmentValueTensor, + roiTensorPtrSrc, + roiType, + layoutParams); + } + else if ((srcDescPtr->dataType == RpptDataType::F32) && (dstDescPtr->dataType == RpptDataType::F32)) + { + color_temperature_f32_f32_host_tensor(reinterpret_cast(static_cast(srcPtr) + srcDescPtr->offsetInBytes), + srcDescPtr, + reinterpret_cast(static_cast(dstPtr) + dstDescPtr->offsetInBytes), + dstDescPtr, + adjustmentValueTensor, + roiTensorPtrSrc, + roiType, + layoutParams); + } + else if ((srcDescPtr->dataType == RpptDataType::I8) && (dstDescPtr->dataType == RpptDataType::I8)) + { + color_temperature_i8_i8_host_tensor(static_cast(srcPtr) + srcDescPtr->offsetInBytes, + srcDescPtr, + static_cast(dstPtr) + dstDescPtr->offsetInBytes, + dstDescPtr, + adjustmentValueTensor, + roiTensorPtrSrc, + roiType, + layoutParams); + } + + return RPP_SUCCESS; +} + /********************************************************************************************************************/ /*********************************************** RPP_GPU_SUPPORT = ON ***********************************************/ /********************************************************************************************************************/ @@ -887,7 +953,7 @@ RppStatus rppt_color_twist_gpu(RppPtr_t srcPtr, #ifdef HIP_COMPILE if (srcDescPtr->c != 3) { - return RPP_ERROR_INVALID_ARGUMENTS; + return RPP_ERROR_INVALID_CHANNELS; } Rpp32u paramIndex = 0; @@ -958,7 +1024,7 @@ RppStatus rppt_color_cast_gpu(RppPtr_t srcPtr, #ifdef HIP_COMPILE if (srcDescPtr->c != 3) { - return RPP_ERROR_INVALID_ARGUMENTS; + return RPP_ERROR_INVALID_CHANNELS; } Rpp32u paramIndex = 0; @@ -1204,4 +1270,71 @@ RppStatus rppt_lut_gpu(RppPtr_t srcPtr, #endif // backend } +/******************** color_temperature ********************/ + +RppStatus rppt_color_temperature_gpu(RppPtr_t srcPtr, + RpptDescPtr srcDescPtr, + RppPtr_t dstPtr, + RpptDescPtr dstDescPtr, + Rpp32s *adjustmentValueTensor, + RpptROIPtr roiTensorPtrSrc, + RpptRoiType roiType, + rppHandle_t rppHandle) +{ +#ifdef HIP_COMPILE + if (srcDescPtr->c != 3) + { + return RPP_ERROR_INVALID_CHANNELS; + } + + Rpp32u paramIndex = 0; + copy_param_int(adjustmentValueTensor, rpp::deref(rppHandle), paramIndex++); + + if ((srcDescPtr->dataType == RpptDataType::U8) && (dstDescPtr->dataType == RpptDataType::U8)) + { + hip_exec_color_temperature_tensor(static_cast(srcPtr) + srcDescPtr->offsetInBytes, + srcDescPtr, + static_cast(dstPtr) + dstDescPtr->offsetInBytes, + dstDescPtr, + roiTensorPtrSrc, + roiType, + rpp::deref(rppHandle)); + } + else if ((srcDescPtr->dataType == RpptDataType::F16) && (dstDescPtr->dataType == RpptDataType::F16)) + { + hip_exec_color_temperature_tensor(reinterpret_cast(static_cast(srcPtr) + srcDescPtr->offsetInBytes), + srcDescPtr, + reinterpret_cast(static_cast(dstPtr) + dstDescPtr->offsetInBytes), + dstDescPtr, + roiTensorPtrSrc, + roiType, + rpp::deref(rppHandle)); + } + else if ((srcDescPtr->dataType == RpptDataType::F32) && (dstDescPtr->dataType == RpptDataType::F32)) + { + hip_exec_color_temperature_tensor(reinterpret_cast(static_cast(srcPtr) + srcDescPtr->offsetInBytes), + srcDescPtr, + reinterpret_cast(static_cast(dstPtr) + dstDescPtr->offsetInBytes), + dstDescPtr, + roiTensorPtrSrc, + roiType, + rpp::deref(rppHandle)); + } + else if ((srcDescPtr->dataType == RpptDataType::I8) && (dstDescPtr->dataType == RpptDataType::I8)) + { + hip_exec_color_temperature_tensor(static_cast(srcPtr) + srcDescPtr->offsetInBytes, + srcDescPtr, + static_cast(dstPtr) + dstDescPtr->offsetInBytes, + dstDescPtr, + roiTensorPtrSrc, + roiType, + rpp::deref(rppHandle)); + } + + return RPP_SUCCESS; +#elif defined(OCL_COMPILE) + return RPP_ERROR_NOT_IMPLEMENTED; +#endif // backend +} + #endif // GPU_SUPPORT diff --git a/utilities/test_suite/HIP/Tensor_hip.cpp b/utilities/test_suite/HIP/Tensor_hip.cpp index 04831ddf4..dddc93c63 100644 --- a/utilities/test_suite/HIP/Tensor_hip.cpp +++ b/utilities/test_suite/HIP/Tensor_hip.cpp @@ -104,7 +104,7 @@ int main(int argc, char **argv) if (layoutType == 2) { - if(testCase == 36 || testCase == 31 || testCase == 86) + if(testCase == 36 || testCase == 31 || testCase == 45 || testCase == 86) { printf("\ncase %d does not exist for PLN1 layout\n", testCase); return -1; @@ -827,6 +827,22 @@ int main(int argc, char **argv) break; } + case 45: + { + testCaseName = "color_temperature"; + + Rpp32s adjustment[batchSize]; + for (i = 0; i < batchSize; i++) + adjustment[i] = 70; + + startWallTime = omp_get_wtime(); + if (inputBitDepth == 0 || inputBitDepth == 1 || inputBitDepth == 2 || inputBitDepth == 5) + rppt_color_temperature_gpu(d_input, srcDescPtr, d_output, dstDescPtr, adjustment, roiTensorPtrSrc, roiTypeSrc, handle); + else + missingFuncFlag = 1; + + break; + } case 49: { testCaseName = "box_filter"; diff --git a/utilities/test_suite/HIP/runTests.py b/utilities/test_suite/HIP/runTests.py index 6150ad97c..bb7290805 100644 --- a/utilities/test_suite/HIP/runTests.py +++ b/utilities/test_suite/HIP/runTests.py @@ -153,7 +153,7 @@ def get_log_file_list(preserveOutput): # Functionality group finder def func_group_finder(case_number): - if case_number < 5 or case_number == 13 or case_number == 36: + if case_number < 5 or case_number == 13 or case_number == 36 or case_number == 45: return "color_augmentations" elif case_number == 8 or case_number == 30 or case_number == 82 or case_number == 83 or case_number == 84: return "effects_augmentations" @@ -692,7 +692,7 @@ def rpp_test_suite_parser_and_validator(): f.close() # print the results of qa tests -supportedCaseList = ['0', '1', '2', '4', '8', '13', '20', '21', '23', '29', '30', '31', '34', '36', '37', '38', '39', '54', '63', '70', '80', '82', '83', '84', '85', '86', '87'] +supportedCaseList = ['0', '1', '2', '4', '8', '13', '20', '21', '23', '29', '30', '31', '34', '36', '37', '38', '39', '45', '54', '63', '70', '80', '82', '83', '84', '85', '86', '87'] nonQACaseList = ['8', '24', '54', '84'] # Add cases present in supportedCaseList, but without QA support if qaMode and testType == 0: diff --git a/utilities/test_suite/HOST/Tensor_host.cpp b/utilities/test_suite/HOST/Tensor_host.cpp index 1e416ed52..f9ae07bbc 100644 --- a/utilities/test_suite/HOST/Tensor_host.cpp +++ b/utilities/test_suite/HOST/Tensor_host.cpp @@ -102,7 +102,7 @@ int main(int argc, char **argv) if (layoutType == 2) { - if(testCase == 36 || testCase == 31 || testCase == 86) + if(testCase == 31 || testCase == 36 || testCase == 45 || testCase == 86) { printf("\ncase %d does not exist for PLN1 layout\n", testCase); return -1; @@ -818,6 +818,23 @@ int main(int argc, char **argv) break; } + case 45: + { + testCaseName = "color_temperature"; + + Rpp8s adjustment[batchSize]; + for (i = 0; i < batchSize; i++) + adjustment[i] = 70; + + startWallTime = omp_get_wtime(); + startCpuTime = clock(); + if (inputBitDepth == 0 || inputBitDepth == 1 || inputBitDepth == 2 || inputBitDepth == 5) + rppt_color_temperature_host(input, srcDescPtr, output, dstDescPtr, adjustment, roiTensorPtrSrc, roiTypeSrc, handle); + else + missingFuncFlag = 1; + + break; + } case 63: { testCaseName = "phase"; diff --git a/utilities/test_suite/HOST/runTests.py b/utilities/test_suite/HOST/runTests.py index e921d27bf..df543a553 100644 --- a/utilities/test_suite/HOST/runTests.py +++ b/utilities/test_suite/HOST/runTests.py @@ -115,7 +115,7 @@ def get_log_file_list(preserveOutput): # Functionality group finder def func_group_finder(case_number): - if case_number < 5 or case_number == 13 or case_number == 36 or case_number == 31: + if case_number < 5 or case_number == 13 or case_number == 36 or case_number == 31 or case_number == 45: return "color_augmentations" elif case_number == 8 or case_number == 30 or case_number == 82 or case_number == 83 or case_number == 84: return "effects_augmentations" @@ -410,7 +410,7 @@ def rpp_test_suite_parser_and_validator(): run_performance_test(loggingFolder, log_file_layout, srcPath1, srcPath2, dstPath, case, numRuns, testType, layout, qaMode, decoderType, batchSize, roiList) # print the results of qa tests -supportedCaseList = ['0', '1', '2', '4', '8', '13', '20', '21', '23', '29', '30', '31', '34', '36', '37', '38', '39', '54', '63', '70', '80', '81', '82', '83', '84', '85', '86', '87'] +supportedCaseList = ['0', '1', '2', '4', '8', '13', '20', '21', '23', '29', '30', '31', '34', '36', '37', '38', '39', '45', '54', '63', '70', '80', '81', '82', '83', '84', '85', '86', '87'] nonQACaseList = ['8', '24', '54', '84'] # Add cases present in supportedCaseList, but without QA support if qaMode and testType == 0: diff --git a/utilities/test_suite/REFERENCE_OUTPUT/color_temperature/color_temperature_u8_Tensor.bin b/utilities/test_suite/REFERENCE_OUTPUT/color_temperature/color_temperature_u8_Tensor.bin new file mode 100644 index 0000000000000000000000000000000000000000..696f0daa56cc6241952f8f69bc76e6960c93f5f8 GIT binary patch literal 205200 zcmeFZbyyZ@8#aDQ*eC`m3Mz<7x7}j5Vj+rRU?M6=cb9Z`cS<7$iU@Ycn!9VQyYITY zuAA5In#cG3zVG*^|9KqFbKJ)q!!R??+~;{;ea(~MC!6UjQxMpv$ggjOzg(67;KmSz znn3vqzuqN2(na1f*>2Lg9(}XD`lYxa>>z^n!gnwH40-O-$jxZN^ABM+49NvZ9 zHrQu}13Tg7h(ISK`69<36=4Ia!uy>{QSQmse_XTRZR^G#+pX`H&OH^Q-4ZpfI(TfU zuSRL$=&Fd3b#dw~38O1R6(D{-6M-y0*=%3gqTqf-0sYGS`&al6s0$rZ8#<^aq+fOD zfXX1b%8;Se;VPwp!z#iRQ{8*T97o7uxb1=W0eJ1l@trtq1E(E0vK0>7V80pL*JFnj zj_gFh5riK@k{b&ACBqlZG0Lscqb{cF-!GeSCVotHphEKr&6*&M%D^#o5!#s^{j$A# z=lRGc9Qn*Sh^GUEL49-mWO96DGrf@IjeI|phssoh$<{>-DD{^;8KF?>KeRbsvn77q z$rI`okxH5Vy@MU$e+b_8aNmi8>u_*24y}R18rUz#uBF(z5H^c(*cz_e;k^fuj)*yi zR1XveBHIJC;r*Lp@1j^M#sn$o2Vk+yRMwN#tl|>Fu^_Ge}gwVb4 z+5*>g*gFf+yAiwv7t;*Z8z6W)woVtg@UTYU9)#{kqN^a4$u7ubHUm%-fyyu@6b&c( zv?r)GMvY)LD}slVh4wB9N3I{z+&_C3;swRNz4CpfQar#?=DHg$yWsDD7&oMPA?v32E&sG+=3nPFbw}#Ze`4tS@Ja_|x^|py zI@1lu8<|tTYgl$-H^S`^?TBPo@KVY1Lvb)FLQxfpnou-ENmm7-Iz+n27rEXj3qyH0 zih~5n6+fSgK$Z(K+!1vQ?t5Wxjh!2?kELukd=DVX83~?4&)P0O@c+mwXy6!{wg6LTky-T7uvGgh`br%x%Bgqk|PRMdc zfiKEKP#J=XU{nU7A`s;PDDW2aEmK?;j-p`J^C%Ae+)0S%BlQ^4Tyf$!d>moF4I5Wt z`CP1?gWc-{@5Rp%p-zZ#LWUa>_u*8?=)0-Y=^H;fdw{Na(2E$l?nMuSsK=Qu9ih|q zbkm(~dD1mL7*4l>=(-18b;vv*(|$su(!XDckHB3u^B9COf0XzNB=TTSI4=MhK7s~F z^G1#+kDtFVWA`HQ7^0lucNqJ&Ve?9?oQ*Yeuznu4Er+8my!Ht;F5C|H^QQgLxS6iI z&<%I$cBRW6bj?rLx#B_{`{;r_o!dzl_fzLVdK^mkLg{fFeVa^AW9e-o-EnUZ)$B|h zQybX3%p0YCs0cuzHwt~g?hUJ4c7sniBk~v$+>zn``P2C7uL#?Q5Lv1=tb9HIg-2v_k0RRzNyiYrA09T?y9hSZv3?S4W@5)YIIYFqg4uL- zH$94=b53;mDBWTbk5R{7I=7F`@28%FbY?4^-%e+@P}^o|-AtDc(-kMW>qqxO=xGvt zn?YA2pQUYhlD)Dk+PKtJ?q~$|*@jwesB{_CZ>F=XVE0nzA?i6wH$3Q$AAQSiPXRqk{@7?&)BfYy1GLWsO=$Jk zu60o_IjmISJgU%9Rj_?OWzgFUdKO2QT&Ty1+P2Wi)l{~KDwa{rDym;c<;$pXgTQ9> zdMaHlXx)p4=)yjM;tP8PfqfiCPZQ`>2K|&quR^L_`-N^q{(SG@gDp;b>z5xI^(<=*y~z~DH_`Mw zlJ2-s$1b|!KqX5kXBMTHQIZj5SWw;!s$BG&p9ZyBQ`TH6T1s^rsA?@WT2skV%A6y3 zK{tHqRxmwJpf`yet^9dv%d>(R`3I!4_9A*4PS_z}D}1-$$SNFNjbm%RdR9IkfHs%$ zXFPPzxa*(uH0%m6?G7?O>#0}0e^mWpt$!O1(2x1_GMyfT(seiLw5R&DRJWeeXVC>a zDw+RpobJCdcIh1cnC#|W+h3ap!bW9#Hw}Y58mbs`sf>+AB349bNCJ-gIR2$-^4udsV6qjJg+XN#DfL<9NEk zX*V~P8alm$YS&Q4tjpV_?(QG(^0;Ep8r;~3?~W;T?~r-wu6pZ`+{G;jP=W6VgpEVP za@^ZL_`UDgr^gg1f5oft$y8-am%Zq2Ae}xyb=#?WE1lj?r}t6qrtcDFW^a_qu^GTB zz1&{$v#0UZU(oC{=H%f~C3}XK*(;aZE0^w8EZwQldPI{>9iney=uQOPVT5Uk%vEVrySGtDtT*Ubv{d z^cwYh{FKL@+ElQL8ttgXjyfEuWsjg~TY0;KO4hzkoSn6ycf#skNo#xM*!2JEY58mb znvajIJ~X0umqO7_#gg61g*%3}9UJ>La~a(amE^jE1zT3Xi_RUQlE<@sw#7KV z9_KgVs*UU$SIxWoNB$K%??1tlehQxa@ThkB43w@Lyc)0qJVqi!3xCF&Q2v4&4#O#9 zE}h&=?K`Pv6P-CgT~5?wM-}U-&gOmA(&SYLoCD7(2w(Wwvk))n2rwwyubgc&IC*XF zw2l3Xwhw(AGm9Ps)9XaR$$1h@cSERoALXum=Rf{qq+#VeM2v-x1~O;i#vvuDx1}1p z*MUYQ)=##6TCJ%=H`g-*P?f8b?XnV&-@@)2s=A>;}Q7XO{M_&@2EjR)0Y zmLhQlvN!hr+*yd{w|nW;9~o09j3XtbCto*>PMmXH(z&Ed$EzhKjs= zE&{aq;P+y`+J@%Y#d)auuLBKl zx$3R&k3$o%v>(>?$9_fp5M@CH%RVMdq4K4FBv1M~eJ-VWw+77zrp-pYO1#`wtHJ4q48gb)R%9!)c zOXKm;;mIa=<2RA+dNHk3xr$0x2ukv#UAwJP#a6Y2y|6$UOJvb{K>4q{#jnB*-i1%7 zUWV%jdc6+T{FpNB&&+vLx|ZsQk`{O&gFh6L);a^mu8P+77PTrBJ$k zSnd|N+^vH|UOpFrGtT-g4w@1k$I+>sLXD}@fx1~3_fh3qinFMifwSx3Hv)9=AeFD7 zi}sW=n^LC!e!}9H6Vv|AT9LC{cDXbbOC!TdrsbgWH4n|EEmD^c3_7(LWsC91P3h0H zDSu|n{w~6Vqx6~YTt+|Sa0=7xPZ4~n6qj3Z7I4L#T;7#o0zf0FViJST+XTcd) z!?ne(gn7>Y8uvja~K4E*0 zO0}Rgi`G^6-cy}gwhH;QOOkyjiGRf!iMHPhu_ZB~HB@Yj!os>}rZ}PpW=*nr>%`?uBF>?o~b?H?|{Q z>wMDqub!6A2B15}_)@gd)p(1miIcjMOs^ywzQ~{QB3b|2FwLI*vRyk-u>d6t@Gg8j z)vl%NxvxCba!rtL0jB{tJOFOP5j+At8-|t|B7Xu3jF31H4_r--Dxq>2GNz$oISMUt zdJF2-qr?XHQ;ctAPQIBx{Z`(jTlr=;3yrQ7=v^%|=*rc(m^H2|TdOB$!j(Lo&z{Ct zf5Fuh!>j40S5nPx=1jewKkaI+`Q;4V+ZlQf6DK^6)4t(7^qTwN3yysr1!??UVEIqx z>_4InGmH?f2H#-_RKp2PWE72l{)ch-sFdAX0P%rf6AWs zDtN-30|N>tp>#4*43MIaVskXj$KR2blw$QRV20O7>==X&M~z2*`riiYy$II6c5Ems z)hEH4k5hENDK@!Pu79|v+#mD_R@PA2!ONsy^ckZDNi`_}tcR!`Q%tNHqOD=lu8nBFd#{Hk^lb?*JAZZkD({X2ij z`&f(j5jx)-8+hMA=8l8(EeDxf4t;y}^?em?^d^4Fy)c~{0b1Y3nLmlpdm5$pC}ixV z!vos4qiG`=Y;eX^=5E&5u1vM_c`7ZbeQT3Yoz$x#y>CnA;I>Tpniy1_Kvit7lgWKP zdm3N;1sz!+5~b)r_f~lBJL265op%Z5kHdBEglS)JQ|LTC?3$O#Wp~Bvo{G2K z6>8QZ!4&Owxa8gccACby#NnsYhc%}Us7sbEjYmZiY7$YOEK?bYig47#$u_18`08o- zYyj>T>R(PB)0wDoGfVH=;@NclsPNX?lgPhX_R$$fy5>Vyynabu{Mb|HUAWP`ql)e8 zP&!NKi2Ke{<)!zC?|er+a#Q^_aNL_1)3=G{_kzYY?2~RkHt>S`uts~Svj=+L^%&l4 zje@y2w;OlDRj$M+or@pR7(bvjV`y{M;KCS`B%m}FmGP*JMNPD9MTAU!{J_ti##es< zz4D@So9Uu0b!?~a_(hCJdJsj|nM&^ZPM|BnbaFqnAEBc451iHS9v*&ronRSb$0NlM z*Y^$lF{4>Y$Uf{&0{R0zLAZsgXj}Fe?j-z^T)q}4Q!b}lA8|8b@6|HzK zQNAb;^+^MpGUT(uP!NN{D3nH{F|}V~;-I#a;pekQef6|_HUL5|+$~?a=}mY2g_k(E zqvH~{WH?YKZ;#T+y+YGz`f@7SNa>3nxs1NBUFyt6T(*(&914GR#F`>uGScUOv#fhm zV6BXt-9wJf#1b{^Fvo@oSf_>kX7HSekR|9iG3IoFVsldehLqm*$$c6#29(62Dp|TL z8Yh#a>y!G`$M-*-tkj)1_Oqw))n7n2gXk`IHga!wT9v5E@T zQqJ=CVdjtBv|k4suoi7xgN6m@+JFdclrBcJA>7sBsEn<{utW)dTLwlS)=2iyi8-#d z%@USl;I>IRU~`{WEn8dS#+;0mYfA22pCZLhX>;1(lSu_!%QkA4v_Dl?QT0YDUP*Nu`K6m5d^L$n{8@88 z#2CKz(de?mBYXKvHodABqGUEAwNbqaR)Y|JNF_W#KO)E|Ce$n`YD!YLMMm_L#6ZIn zN5^<844BQeb48NxmjV*k*V?7)A;Hy z5bFEg4nozKN9%TK+Dt8*scacFZK3={ls=uRR{RmE^T!GOf8x!5i7@;%(%{iig@^kF zw=PA*7{N+*ZIzDG!IeG3f+iuud2}LIvVsg#Bh50R%y~Z}!tBVFKK4tI<*8O2-oKo! zVKAzqa3x=>GkqL~tJjJrUdT}E%F%d!cJ}G)F_#J_w5N!D9iP1}>HJ=y+SqJM93gUj zux2eaZxZ_dE003F#K_ zQ$>RnLM-4v3&E?9u%~~Kr+Qtues!>ZsgKsBl$m>{p*BpGe>tZoU*&eW_T6&R>xD*V zQ#5W?7~QEby;)`2U7_Dqso#;S-BYah+0*#yFQ67{YTGJNr`+1Tn#xyE`En{-N`;nO zGN$Yqzl4tG!2iPu6UtifQ=r~U&oOWPG+E7ZGN*YRZaWTq>aG^3fed5RFNT*Yk|!g7 z73%DJSMBW^vl<0^`k(S1Q*~r$qvz05k^NhvWx8|JI#Ng6E;qZBGx2QtsLMqY?pIB@ zUTWHyI^sgH_U*b!=d;Cn$7iogIw^TEoVnv>fZ9#M3;xws!rR1UON3V!V~x0o{CAA$ z-`oVAGW$p03EUI@XQWa4N<4Q~dKWzI<{`OfZW<+A_qRfZA#&&7)K;Vzp?(?4m!ZW@ zs%ekRRez0dVomdRAZ<5#vQ%n9`+97^qo(OMtBec%&>TDTUYW(UJcFvRfoHSSZdaL9 zgoyn)pU-37JStv9_3No>4b@l)JXWj`EMvxu--E|df*Di!SG3Wu;d+!k|JPtWN}Teo zi^|#6s9AvW`FQNA`g5q!kA6B8GtjymC5upOiELwB;4pKwQ1B{Th09(F)w`sgXIa#G zN?k}%x|BV(E^26Jp=N7_N}>zm9r=gp#(mc`w>q@%*-RBSmuFMN`0?|3{M7tbuteal zfUV&Iq4Zz9jxuMmoBAG~4(=rha2kzw=2dk87tqI5~-b zBa460E;j@P!RRTT@U&%GV}g94KmT~(pt}`@pFNGQ{sKV;d2onPzLrWB3BFDK0?M6F zDN`tU(qG{d`5n~{aoPd5(H=Y%r$+x@>XD7FXxYKfQJrTCc;ezJ_ZOhMW`tvjD-<91&Q;e6BbCtEtHC$gFtim>cdM5-m36Y zfR_^dN5XFuJXPQ-58vU4P(tB&oH9e_WOPr%c?+C1!qSPlHeYZ?{r zq1>H!{pWYNTD(l(@K^O-YCBG+J*m;@ZRVC6ezR`-TV8RQQn6M((G+QhD7C=JDL6Z` z*TaRw->y}Ey>ZNg)r#$Nq_RzrG9AV1Wt(>^mu?;qHybA=BiRC}1}HQ{y&2lf(Po4e zBb?Mn`9zdXM8$YiPC%758V%4j3zrx7YF{K>FbmOB5j+clv*A4pu4Xu_i~SRY?NNOk z)rI3kVSLG-1S5p#z)c%Y+HjiyZ+-Y1BiI<>`U1Za9(#S_F^Z>Y4sTs}X~JuSK%%!Y z0!ATVq_E>Q3?73KrGjEz0k=*I^h^`DyI>&@dUG}|@t?Nh-M?%elZ+5N z0j`=jrUlpWg17}5BghD0dJ-NdBWWJu7a(Q-}YJfqGvHH=#(qH zNTWCT)Zz1Miq!+pDV_VYFCWysWIw)b^YHA+$TmjHe6%e<+bnd>lDc8p|K{QWU6!(S zQ;=-i#?7jT6^2ZC?N?n1+CVgC; zgWF4_I_Kl$Oyo^N=1gSHM*2+QCcM1kbP=nEMBXt(tN~*5zhrR~v{uA;xR1ecO}Meb zoQOa@1nD7c0wQz}ZHzcxs522U9l;W%>aGb_b-1g-UBY9~Chyh)Ez?S`vZ?j>180-Y z9U~i8%QdeZ)@n7ZZmDdJ1qzK(GZnSdQDKH^)|Mt{nuNMZC^1HsK5|S^G!s=zq$`(5 z=gvTqA(AE_bu21$(PWHM5*`~2QK>I1u-x$|V!f*WUmmZ|!@Xrv-Sg2j9fcMGl?5|U z!26R>V2XSr2Z zAp#8HqXW-za2pMmQE*X#+b4M(iC{JONqBT0iV!8_Y2&PkAc>b*qnhKaDY|Fl#u7Z* zs(8&-i5m9Nb$@zvf?j9PlLUdjZ`0}fOlmmPxqD2*ivFzPil_Z|`46Oo{a)G?^g6VyBN*uvBs3EbsR{4b9z zn`|XnY22ERhs&idTcUX;a!in6flLczPeQ&q@=TC#j64Z_8T=K7!uXOu`MOBdLa-+M znL#c1Yr=0V0!D)cEtGAaI($aMQxhzP0)y&s9tmePfk$^`L9a52AtT_gEXX6fIuWWU znt*nFTr|c-Gn_NU8CI}MKqoAz;b7Z3xsNF;|0>u>JzjL)g)Vy0>n!>qm!2hEb}??Y z9{$41qIkAIO4bylnB#Xi9l6huwUw`3MVxjRl|`Ha8$(+ z6&z*nR0$r;qw@dd(PfBWAKA|-&_b&Yy3BEIGM;Q2@WTO(&XuxMv6Jpa(wXCQJ&;Zx zrTdZ8eDKGt^|u3NT=SlM*~6@5ug3S$^XOITL%+$%vym_h@zVuUl`|a$7AP`D_7tQ| zLDF<2&IJ3V@pF(k4~cV-G7HHwkuVLhlM!u>C_}_dK=L@jJ~l9kY#jd|9$6&W)v1_( z3N69;xw!z3*Y&xwT&8w5QrK@dLy)oH9t7&cX99dD!ruTPCclWOD8W?$M~35|0`@Axp0~WOh{K9-Q55W>pAv#qg~59DxIDauB6K+N$D&nF za8u9C1!b+GVw(pCwW!&F>i5&NVCwdz*O~Nf(%;p)KNM~KquBOE%;MI4qgxMXT=zAh zH>q?zAYmaQW(XEJb}CXPA=d&q=18B4_-Tloj_?@>osF=0I58h5<{)aOFv6!HcoKq5 z5o{pXf>=%Pyy}d=ma&0%3ju+rij2R2a5HU_VSa z*;Nr93h+}vkcu$a&tZA=kVm8<3N+E6i#7|iEdrHq{t&z9#}fbO{IHDuDj~<-#;cNgmBe)vE#X#78$scwkLdW9ZaO@Z+ z(8Ob#ydZkVl;ET+IX(#1a@aUn7+VHmhn(;)7X^3>69RuGk<~AIILC*=Z3se#A#04_ zSDcy&N|-~J-RRU2fxg=zblHd6k5RWDeV;+M!USgD7t%ZSbCRgrz4eG@-9Zi3x?R2& zt|o9d6$Bx`NXSJb8X%3`X4ZQa2%3a|sqmjBjNr+_@RgweXA>NqC`7?N;}EHh z0z*`ppuq@rdIFD?dIFCbW00?fB5f3n6%O87g8NnjF0SZZF$a;R!Xoi7fvd4ZB#q$m zzj549Ae3qSl0TV-2-U(NCDBnZ2C3i5~hY9gxx)!Rk84$O!PsAbw%oI5FW$UJcaBaaasy6*EkL;b6aKyp> zsy1@9P%sX8qftFk;PK%Gx$f2dYAglI8#Wao9A3>p&{PCYf&V1m6zdUXng3Uv)hfU;AxGM_M=Q$KXlEe<1QFj#~ ze6C;GyJ#|g@iwFS-PFnk(U*DxsP!ma@)G!}w-{H9O8IrwaXzGRm8Jh>CSjnYX4a!+Bm%UOZXFduRq0(BuVuw=sF13=Ks>OCHB|B6q zw~eUSs#>;vME2%kNvr#xSS%H}0Fg5gWrhIpt(s&`fw)eCa%Ls2eq2a7XPWRPIs5 zE=BBD7D#j%0S5)_93-eyW{`Q@G6Xx>J}L=8hs$7ia|o#{IIg}LNU#7^?x)Lv7k7^M zJ#r2;?WJlvYT8eo?sV!X63RI)_=TsPv`e-PDYR3_-##qY zdPv6lftl8FY1RXi)(=Rr8j!bnQ0^L;tcA#1ips6>m3GP%yGJsQg*%js?Uc*5tCZO) z7jIR}wH}(hwqMLLsR&CUx5kcZln!`()=WmdsU&;DfrcJRIO)oHga7eZt&3aB@np;3 zt5$OLi+iQBFg<92<$fYG4wXO_s_7 zDmXq8t{iJ9VYfVN2E%$NHVqZFc0&a(i2YSh7D;w^IIYS#ff2Z{cN|?0qdVbG0w=$Z zo7J3;P@7 zvNjG%T-_&jxirgT^jySD5mM#(7O0wx#!09$K)J49D>(?w;gk-0JKD$}gBm^DULo*! zZT%qTF=0A_Cc}@d?Ns znH%JDtmO08_Agl_TR0CqidXi|TiGXbMc>qweNxx*P}B>>2Wljb$aGF?&Lyzh?`42Zod?+<5;8APNm#_ROV*+bgKcWYx||F>6g5^ zZ^DW`8Eg9IuI^j9Qnq9ccoZ#@$yzF#xx_)5B#zD+`;#%3b zRZ`(>ALpXPN{-r&JUpP?JP(~qrMj2&sxd(+$B#^+u8?JApQlh$@MG^TmwLKo=#7o? zP0M6BnGnMA$X*Tmz!pL{1T;DqAkJveMau1z$rW?W3;q-%>Zu3%H+NIdOe`NDf&Bh~I)d$Ap*{P;(9Gbkk zf6B^!DJ$d>SM*O<**9^eZ2C%>g5@$yV!<@Q4~dt={xJ)CrL7#8w_$jxom!raV(u1& zT$^Dx!l%_892L7vsLPbDAM~FSi<%erZkdg93vq9=JQpH4mgLMpkr9gYP^u%yBUdEv zE${dJHr3nK!&rDT=7S5YQF9T&tg+Y1q0Mx`M6&zJ&QItJc$t35;wa>VUPZSS=s3xDQV-U^2Won7Gb(0PpSWexNy}y97E4DjmSU$Rez|nU%09VEdT}{}^Es)Etp!5a@t6*dINo#B z@@Dw6l+`6WMx1jurYm05cJOK9; zv0RS)?fr^(%jfS@II)alMug4935Jo?<03>^B5od%XCQBqP@IbAJg1RR-t{y=r?Y{O zNqCvYDZch2daVb?UGy;dA$#3lRl9yjTl+S4$v1AsHS=&`?SSUx;MlOiQpmSAF6n*A zTItWgIjtPX=%dOMw^mDiYd7rPwqeaHqy>4jMD)^LCzkeNzn7yAW-x5le>wHD5RwC5 z@`q94Jb%O_M9cwOrQ9t_RlC&d>@_O4C>5_8+`d=)w*Q=~LGwC&rZ+npV2?PTbHYfdOmW@+1gGhA z;LlEo4z79`(nS}#5lB6L9Q1h`3dz;Z<9b30_gRvVXQBck9sRWm3u9+A)`xmQ>6NMDkimPrS)_16+Y4>$H&mEfs{Y%Px0BlU zQ~Acubuv_7MK}EEg44NOLw}B+Or5A2=YXiO5^JD&u!xOP*4D25eopGdAdpdo9x;?1N zlkUdQPbG9ej?F0Dlqlhr!_?zV_d_XT<;gjyHo?8s17B@d`^HwKc}Xv>ktNJW%BtQ8 zt7OCGBY^!Jb76#V|J4-uNc6*(ERaQ0kSa+*N6tk2lHNHRhE;NB#h$SxR`SJb2A$fi z-QzlmyP+<)nm6s$uC*OmvUV^hOc-J2&RGFI>O$!`ggfroSu#fDronGRr(WMN?CHUg z6hDva9dvRV)!AILkrQfIJw9~R>(Zg&f2Lbd%W^uulG^9crMXn4L9e#AI`lrKi+$s; zcPtK%hTB*~=?dlD+qU?RH*Ow4hdoNx;z_vX4bM@xJhZ9O_OG-Rl((MBw$MpCs@X~{ zcGSL?>UOZ-MVI~QVGO;^7V4fIZgef^KPl^|CzyUJr3VR}YloES`13}P8z6QLQdj zX;RY!I&V&wETd+_l}knXaOD03W0Wq&6G!~&1u6!$g3icN>j27DOqz zw*li{2256a9)Zj$zLStU} z6qM+5l6=UkUF!F@sWdI?oiJ6%lXFroWWHday)6)EA(T6_R!CQE9oD*klmi-m~9}|{R?Jg?X_}b6-)JoLP!{zmTesmstWdkm*M9)gRIi^78Z0WT> zwOdj5M#|HoT0LsFppKbz&xS5K!OH73<9OnUqvL@G~r+}ZQ(RRpH=do8^ zN8k6?{*rmjpC)7@xeGUWv26CbAtl>J)a)8lwM`{&Wxvct(&g)ibT~}7;Xn0$*qpx# zZ8-J)TiS}((X&5hE&nZj2^DRk`aN{bk6I74?NFq)19ZjXkC^H2e048aNnNq()3Fez zIIdib^9yio1E^#^RW7EJtEh7eW$RK2yPZp^#gfWq)6Knf!I~vp>9{#&^5nHA_>C;Qz2wJU3TU0lktGmTwS>N-TtJLo2NKYCH+^0#ruui{KT7S5;h zHq^S9x-6+%kJ9mnKd4YbU#>(LSst0zPs~VjrtX9E#D~so_%}hX zdCO4%ypF7C$M@tw+E zji-T7P4?D7gb8wI#7iIF5p(aOz4BEUJOds zpq@39I{w+=p;Tr?7mmI6Hafptnle{Y=_ab#_HWKAYS~NQgwR#jxAC*y#83MnX%bzs z6|OpKNL`97Zp8GlFtRQ^DG@{7XL;no^zoJnE#)mIiw07$9#zbzc3ZmcL}xZXbs1j3 zjnihR+tiQF9TD1Px$)qf8{LZ&EaMqZx)MT<)5y*8ng;H2ptQq;GWXJH|KHOnA2jraTCp@}G>Qg7aCikxuWT3;V9QYH>}l(+M@!b+gkuk%~rBi3!!Nrw)hzM3|*bMhbUAPQ^3- zNp#MM>u7YzkBYZb#U3izA*3MK)#(VM&HB$2@OB8^D1#Pjk@xE7He)l!qk1am@jzh` zVdpR6!D`_Xk8NKxkDR(p*zs&4`(eaz#S z+V(?=<@X^*AJb-E*w;Uv+efG2(IFO=5tP=8ilphZl3=_T{|3-0P@6nus8T5>$v0EO zE;_l3^43x%%%#txCcRCbPG!reX5EKq13I~o8m3Xb4i!ra^sxxAK}-R~O12DU zIw%Pg4hmvmWvFK%b+4vYb1EN01?tphNl*N#!TzzQ&J#~PZZ{3l#@%B&RI!839;fpj z6tRdx=f6CvbARgyid;*_7rxQ%^_?8v^}){r@k9=HRM5eR;@P01(i9>2x-O17I#GBfCw+lT;i^I9YvuEo$fV82Lr>%Dj@m6Y@_%G5r?L%K52^B- z6je)v7av-e;hl>_K;qc&Np^VMGSrMedn2p{^JRZJj-Py#U-^&wEp9qhYV-D zW7-lrxm6JAZhI=7`!RepB~PHHY1D2e2xH4Y0k>?47&3jKplFFy^0!2Sf+a_aMp8W| zb5>ExsNW*`Q|c({u%eq@^eBot4$yVi7OMe4>O!OP`>B$1V$EE!nJzfcMF%=#OQ+XS{q*112Bi+5R4FQy zNFoa(TSj(JnM9Td7Q;w9*^1kn@T)6+^T40p{4WJ+SW4N`Pp(AyB3$0jdUyfVZJ{D7 z%3Ji@Z47&EIczB;jjyp|gyd4K0pE5zJxZg8`DAT$a{z8j;|_4I7w${pp$r~N;gu|Y z9DtANtY#(CVG?;#36jTPKZkXD#TU&ZHzD&&8DYu@S3#&MA~b|P|1fogsv}BWsJgaI z!;9T`c^I!<(PM+o_4w#3k--oNYb-JW_=m#;KYToje>fb7!ZRCOTgg7BCKau2x9`bayM@7r2d?ht*qNdGMHkZnbD5)RCNEQQA#u|)S<2Arr-scsPDj}4eQg^)EE%5u^ zS#p0qr(rggSf3R&^ z$)KH+9}mNA8QcV}_rfKh8@S0-_QI>)_-P2p1mq;q1FXZ@&*8^lqs-fsFPg`IVF*`7 zloFy8aY6~PBM_yGxDkj}McW*Z$N!Wk1J3aBmS@M@u?#Ygtmv3HUbj5c9NgZBTU+sP zCz__#i%1WcsTWmLVAN)}M{VoKJhlra<~E!^Rg?&4+7(#-_%{XY!` zkf~(xiX9 z!;i;SgH^g@(SFH93X-UV_>o8&i8v*s@s1*r6@`~T$xX5jS;eux!6N}bAHmND@E0#v z=8k8_Y_cNdwR%PfuUHw@&lfT|{LX358rDTKsdfn^>Av(FO6P6q!ggw3Pp6js6QdwF zgzU61R{w`bhKrAUS{uAPSY0!ZtYg{iGL>w0+4*5MSz7rZb9!ntWhg#AF5Ppe-#?i% zsok2Y=TJ~@N>it^w)8NN-p10oUDUpXx*X`7=d(2jXZ?%9amS z$WxOUH7%#s&2(coRhUt#JXPz`b$hyVoNjQPWiHh%rMeB&>CJBH{rQ~y37`_U`xDSA z19n?@G|8Y*3e9~{&YVi9ywVHb4Fu^6`dnB}tdaKOqYY^PmrNs)lo79tL{%iKAzcj_ zYUrFNQ7OzIL&t!B0+40LPa-d#!UTUhjDLB8#gagSrK47G%LPtZWLQIH_7*T_31ZAy ztoSAw{V08q01N+@DP!VzZLsvQ81lhSTfYBM5F^GU@_#?^$XmXI@4$n3=QCM0`Fr{1 zC7f|Jr1Kl7X$^I+r3Mo!8b$3>={Be7>?m24!unHzIkj2S*#i_X{g&lG>Q1Kfao2Tm zsxKNO^2j@FeNZcnLRpl^pc%L(iXY`oRIQ;2_pO4`W5B%nW53cyY=XgoZ`DDWQOa`1UVR_+eYx`4! zobV@04I4F&Q(pQlgZ=d!{B;&Zn&D@3#%SHJxJzFA83F zb|g;qmRLqvG_z8bLQOB^N+Dkw=jCuq4*yXFO@^RL*>La!ACh=JK3_DC{Du@8$9P5L zjm7I-5*cCqKY`DP<>1dyy!6D4eJGiSx}|7ZgMVQ zz$5Q7^2`^@RMv1x>@85qqusJs&RK740h8+-?rTn5ttzwDH- z93VNLS>|WsoBQJ3PFz_a(Z;b-!Vz8-yh!-AyiQoI;|5WVK1GbABSRk;;y247mqw$$ zFD~{+r!>w;fs<|R5-V9Og(4Z$0q14$j|u3YEPttBKKQ{{f$$Xkzi1xgMj=KS=_4gU z7K<1|R3s@Y6iOV7cwDeU*=od1hu1{mRVMDQ_?|_hxz~S{8)N=S?=crV7(E_LZiTcU zm<%Zc%UqPGQ^QoMHG6yjJqv{Ql+T;vccz9>;;R@az6qb{^S^~uE+O%gN@rED9ak6P z2Rn%pV0QRR`61JBZ9ZD`(V~m4sS@$wl@NfBF8G)4!}}@HbaozfZKhoF_sei$2s-3& zW(ew87^QHo7tZy@Sy{p1X_iHWEDG6NattyMZ#6)Bdee50kYYoalhSYnJO8jZboMNsKgiuB$)f*>G@T~VC`Ard7$h?!}8evmo~1b#o?2YyK;zFLQxG5Co^E?t=J zNxrNi`Rh^mbPX=ugtH@MvEC~S@ak$jyc~~jz~{Tz4>Qs7inrynOv`g&?@q(RPB=r? zbHYO%@r1Kfh4(kZB}yxtcfipWI00#WeA@wEcEBgy@iX6zjYhe_w<20zw#NLguq!J2 z$<8Dab#T-+i<8GPPWDAcKCF*QB;Q27GL}Z3Yn51Ig%!Mkt6@G zP<*@=|L`$h34-R+Q-(0q$;7E)I6M$lgHbgCxyHKqlu_6>gXPIA%M}mH$Ga_$N56kZ zt2YNpXnGtt-Aq~>thp0C-Hjh4LCOK6Hl&{# z-y?ploROgEk7uPJ`HL@(4nj4TH;zR4P~;CZ_Nf?!98NP1>0`O#%ej%i`i;3V5U&ly ztG#fZ5L@B#mUyleo@PHlbDW28q`c&WAKT-P&iJ-1K5m5{d*Ihw@aqcvK$@(V8#hOG z;1u=7(J}Z}ltG6m3-R<~lng|6AEaQX!*sh zcxoyhnv16vkY@+Xq$Mk?M%OJ@6>^HR_K0T<}mQyfzMh zx)ZsXxwgb=6SmMPv^~QBV0_ zoRS73m4E*r6pg^1A;{|j2AjD+WpgP9PO}dOc+n9rI^gBDc-m27)E8Pv<@o33 zc*_NEvbODyOT$q)5c_9HgpkOkOamG<$Jj2itT>@74o${=_ei;PW;*_vjI(!&`<$47 zPgzdVS!g;Uqtj8LM^Ha)2vbyCO7p7nR>eI0NElw7jx$|xxE)Gcqr3x(T~OK)rv~6$ zA6)K;$9v%Ep?GCFo|%nLyzt^Wys%N`$bHLkjkFe@2N_ggx?h{BJV z?LLdeJO20~!YEWpkX0~I^3O=d3+qsQkF1M4KMTjlO8oO{uxxXb-CLH)Ck#YFKco*t zN>_oH-U|sXNNa^WpaM7qR0B0esKZG&9PfDD0e|U;*L&cJfw;hp5V>5p4Z^t{4JTaP zj-WBvy$Ek|$$BIHx1$|Ao#4?1{!SRykIOc$JttHxd(Mej1 zbi;SL;*~ylau`mHMq*ERbimfG@Z=KRNwQ5x=s4uxi|3LX^?W%m7kIaaS3CH%!~Ti* zcbef@tm#;CvZ$o5(ZyIvFd6Bdo$G3Aj83r=}ab zCUW$-i{Sfi6JEF-pRSj%_KRisel?yMZ-hoM(vA7G7v9~34_%FyXAk~U+V;e1S&8{v z0RD4e%mbT|z|A&S;j7f-vAuxiPIhMR5uy&-~iTJ&Um{MUhad3xv$qi?Cb=u?%2{9Te`!0pzOio zIRswA5Og!j*Euv|cZ5$Tcsj$sz5JJ-U1J29bSoN%HX#>!8AE--B0F}Q?|{cwDkz>-c&+CK?@1mZt__;w?{-DI>@ zBJjBf9_Gr^yYaU$ak@7GWUZxpZxqkPS4CsaZ-ie@WKTl=bcByY)KI*!-S{JkFAU#! z;X6O$r;p=@Oq9()`EA(E-5O`fimfNMcla!E)b5e8sSY;>Aj%W%aF#J@Cp_O>s_6$< zpSr-`MRs}G-V+asB@b7YAgCK+hauP4 zgYjE-#IX>LGn_X<%2Kxee(^DGnHYr+HsK>T+iwG--F)UF!L;R&NI5{?$mn z896g$#rIBjlP)q!^;o>S8jsGxi;Io-?F@d-LB-uTzX2}<_xxw&J)fkG`K55&=P^AG zO~U6PJ>T)_#;tV@jg^f%Dmcp16$Raq&>2C_QW_@?kw2IVj<$38{AdK+f*nKQ+8=AW zVPo&c?2b|wa5Jj_Mjgs5hAaKF3n#dwl-noXEZe(DB}jsNL?@*ivSvJM6s`0=qe%Sw zHvCFP^ay$!>q^EvMfk`aY5m~U5rN#4^>%z6ivNU40r>AweCLfHIUXH>e+1)q4xohK zpTtV9=NHoD&$xV(j;slA?S@U=5j+%+d*Q=Kyy1sa%VlwE5LdVlLg9G)ZWQBm{FhhA zbRemqA8G%*2QImCZ5=+y8fQ7a;GOv1e+h5(O_I|$v2EY=lmAyPcgJOC;}nE)MCjhdrp2l#YA`~YJrKJA}{-;<=s zD&nqR+z@OAUfhB|=q%(&qp_r=o%b0uPZ?DwoA5iuD@9mfG8AE`lZX$uB5RO%`o^~K z=qO={Jc^be_S+qJaUtH_h+jDn7LGsS@ohMsVWVR!Uh>3e{A@f9%tflPU)&D14e!G9 zUbwOa53I+FemJ`pp(EfoST=AzMjzxKM&D;q&!FpqqourM(QSA=s_%QLW1jTy_QK9y z`{&@50DP5#&%*IuAYS&uu?4bYWcqCg9gK{r*xCyoy=C*~cLVVO>*#Ix)CYg@M&dY%*7j-a9KAbAL+M3NJi8eWEx}ptVm%gN0}#{)MH6Itl@zMwwP6g)QgURW zN}#kJLmk#oEKwQg(xIs5;=xHs>5QG7v6UOsIbxf$Y!D|;88Yp{yX$b8+i6b0=kE9= z0Ds?wH@Bm58fuo|$i3LZ*;Dh7bc?J#OfVL&om+`V)+2v{Y|Hst5T5msHG1NbBV|j* z`&bmm;Ok`kJs;m^;`A~UPG>nh@NX4!{~SLuag3~*{5%!Uy5pf`cy$N<5`yCkkv|hL zw;*sRya&U3B)0XI{@Bk$@x%gLoQ4m!;NR&en1sM#*ggd7`iN{C(AYV$I*?kjGq$yX zXIpr-gKKNJJIn39iG~}JBMVgu{gGuXuSX4(jCKf=RVqtYGGg6Fj`?_b4&vG&pgr7N zu!$|2j(B-vt@Fm=D-WEy1yA0MA6PR+;1f5L3`hD)-1#>LID`P*jhnZ_;shs%nLy>oGV6)tQ=(HvQG$CmqZyN16VH|p;d zbMA9($5Q3}D7+PbC)VMim3S`{XO~O7R4_~Yn$Eawka%~+G<>oX)trt#7_Y99|Ml4w zvbTfhaIEQzRlVUltWk6H;g)gj#bvfSVQX8sw!wxLa3@Mfc(=o|n^Is`i{UG) zFor!Wfk}|gNZ(|ErrIV%%(h+_-dc$y`eQo@gWcL7tQ)>&*~?;?MQb8{3&fB9_=)jt zh?L1cd7^xP^gpBwls#a%%8+}4k>FPD%ft;uhvM`aycsSuUk$44;rrY+mV@0bmHmz{lC2}TZOD${Sk((FdSGd{M(vM&&hTvqZ+7W0AfY%* zbJE)xe(e#|0ioOsuqXD8$2-CJKGDHa+=-I`81hwdJCxbDpT$%#9$Sc?8Hkb~`(1PK zmWTAEaMQUYZV1^LoFALn8D*XE_8qWz!s2eszzN0QR^hqP`1np-=!%nVP}K$%olwkn z+(;R)jx|cowcJFjn?$YtgAhMXcF^@7fo+4ajm95{tvzHLfymL|{>j{}<>>b2{)xn2O4Qv^`g15gYzucpP$z_RMku$* z?T5qxNF0dEYw=CS$VXP;*lZlU8#$AZ!JWW)9t4-7&%`S`@g-}8g3F!4`XQhn zcJ@VRJH)j{o-<0@;z(;e*ac7a#OovR`+EH0hCfZlNk^P@Mva3E)1PXMgB+&67kg(R zVJgCmj+9M3u)Y&EvEuFx_W_3E^vC+n*w|GzyIa! zb0VHvf~S|@(M`?IFJj!{{N^nDy1K6=Xwr-8+edR7c6Jf0Yr!pgwsMnF_NCLoZ^o{X zNS`WOka1TB`upYq@Min*CTt#x_5HAAaHHnv-34rVc(}l)Jp#KJdmZyE0B$+X@q(eq z7>>O0sG5t{!v>sRisDJgmZ)E!&y4$coqQLzp;i7W^s4uQP?v{4h-5o z4JVefpR?zg71+()rRL%F$e~XKbmzu0+|WC8lx)A2G6lQ1@68DKvN&=)*H~t@ z!^sx7+)CzbbF}Dub5u9OX$Rcj2KTp?xiOpo%oS;8+VP;4=Fe^ECfkN^TGtYB$-vQA z&!Iz}V70aTIlsQ#zuL7o)_27!ZkFB;?nAI;09JOwD!z(ikhdI#D-nG=a^~YtDPtdY z?QnQD&aT8$>sx)bwa2H6oGpQ4ADJoJ8?q{l94T(dKXvC&*-_S&8+!BIy9v>EV*5>S z9Rk;(*vOuo(QqGy)x8nSQ(8Ed(1`!BC&!Cz4Gvv#AI^SqP6zih;7Li@qh^5foJxlbukKJVM}>)c5%hWB8Y5yCxKZbk|B=;Kaev*p04vVN%PjVF8K z5hq+UX8Cgd_Jb|N?=G~ENll#5bgnrrah_o_JlqzKw!!1QaA6dVa9{sSO5mN*L8PzySyz0+=IQz?qIgB(icrMZyIlnAxjOcl{ za+$XC52bg#6gH5XX|RFJ2HlHZE|=Hi412MrNryIfx2U`uFK)owOYrJ2yx9$}w84|W zW5DB%c*p_w8LoHE0X2>|)&hrFC$_?w7BXY)X$QR44G;E`8DISV>Kz?xJi2o86W>wT z%)?)1;DB3)qQ!FVTkvSO_K_gRgAO$c-Z#T_FgCG;!ZU#GMjZDvyAzCBdAog(Y*U%W zUCQS*`_^ytXWm1&e)Gsm1Pl~#EJD2o%9gVcx4>^Cw)B;L8*+@~E@ZQ@cV)Av+u+8H zE~jA2Agu3$fN}7>xv_KfG@Ae9$U-%SN9q_Hc^py#+h~0d-Cxcu<2LNam$l^14?L$V ze+n)v!qZ zC&7K7bhAXy5R`-%G7&p(Mao?Wn}Eov*fJ37dSfHYOx6pxaiYl7KO(05)o=InPbMDSZMWIFvI;qQUBgcwm(*Ek8XE z7nb0KE%+%ydcQwjgTGA1`(5x+OOf~6;>9*Hw~)aM%TmtuJ<=Tenxmu{_A&&b&WUP_ zb|4eHIS%-2e>p#H|EiX!w|CB*Bd)=CgN5g^E>a2Z7!2Q$(pc~qEW=5GH_7dG>LU?4 zSu(SNqYJyyTZbWYK8m@A&xQ_`m?{6-I^-A6VUK%se>i+l;l1!NI5OwPPMvM55mpm+58}mO#s_stM))owS!g|Q2e{5**C=*UOix?pL zA&2&teMkr~kt5^mVCgGcvp(vEC*4~VOvGcZ?Vj9-kJ-N(hab1$)jRR*V0_dO?*m^t z;sZy#;Eekn!C9kpL(UOB+(O*26xa)tHN$?Et}SJj#g%4wzY{pcw3tWxZFHhba&zN~ zwGO!pv5j?NZ>;W)6yza3a?A$J)NqZ?k&t z#ECIPd4jhNf|XL!5@%cDvV%+sW=;BjSG>U)#gkFGKz1SKUaX8-OINpunSsq5$?YQ| z99K3ww!M?T;f$=jmi5A#A@Je33^U+89A&HULfS3wl-^-^$d%3cZ$gLt;yd_1fusLO zobsqw_q4kZFbZBghjf~Fu|I3QaR{7*4Fjde#+}Z7mx%98*v?XS2!eSq>qPiZg5Lyq zHtum`o`tlfEjF=7&RI@G^5_KL9@xpvtJ!?$0gujdb^u3kqDIIDj6wZnxaW!BF8dZB zVkA7+ufpT!SXK4Gl?C|ME?G_U#Z-LP9>27}*DdfGWR~wi;%ke;TtLBEvX!voWMH;6 z9&UxlTH+~AHEfCZ9Pnvpygn4ixzF2lgpQQK;|OyJ};rzn@xhcs_m%8E#3N{z#mi_>29gb*@gUl4f42C18!;N9>t?(W%;yR)+7pron)#f(f z+>>>VSnnjVxjo#wz?}_ko`^(O<#AVi5I6vSec@*`xKvY-U1U3kNFR#_`#pNemNlxkc5;%8QQ4!Fe7v>DE|m8ngqn@dGWf?Pd%r4^oNh3A^%P2hKD{JR~# z;`R#zaeA5zVQ_OWe};?PnCMOvEJ5Zx@saS`kT6RM(gz|&q|C88sn>9$e}L}E_9lB= zr=fB~t71++c5{AY89VT2{p;hah~KG(fg7&>D$N!{OPu%T;#D(I3~g!pc@y-Wp3=V|g2_bi%5(SlJpY zT4Ei$(0Ni45AvYULwRl>338WIZm`MVpN+WT@Ea`yMyW&akt;sG9p81t-to)2>Fs4ts|vNha9($MZlCs?T^0fy>u3TB(fzASlC)*fddvfU~vm9ZH^VK1m$Mp<~b}K zrMHv4njHA^;2mpZV31 zdmoM)&ayXpns6k<&?yL-EFWZqLyp@=OP6@eY((54Wh`&Ej>4*b*f3nm?QM;_eq?il zopKv$IWBU<0!Peuz`d<-Zwmt@UDW|=+F_lGxGH5ePuS@p4Q+Dt7=>+&3fP}~GuYxg zu^5-9<4+?b`24gD-fD(Nn#+Q<$J^miPQQ1Q)uXSqlIiy!0$webv~5%9cO zqRKTqpNl7tFmN4&CwH{}SL*oR(#F4^I_k6iv)-+|^U;(sd;!k-_1(#H$Jn^I9RZW1 za%7yt`jP%fpnkU^bpbZA(~l!oOdO1jBjC!AwQ;vQIJ)7>6CPQXw#4FAhBG>19wD~E zf|gj+0*hK=87a5LvbI>-28&x8uIh*t^x`g9*8?khV_6>=X$=|+9=pJ^;doZ&GehxS zSNz#ge2_f=7h6hH!(|lD0dKQOLN{yH9b0#}aFvx|nq z^z6Oi_!P$n+2Xqep|>JqjtJY2?8^2TDd{W&=|r@D(me5?RXwq>r?ftNM#v!!k8XE- zD`o(9K7Z%H9WNA4IvLQD0nEqy7N^}KwJrTIGDIr=9l}T!WgvL#0FzGtn7*<^nH#Q4UiI*(I%$^hflyZ?v8$o+-~lJRXwBw zjoo)F4CzWsSorsn%uEU8an)<_YRcd%34Nd6ed}`tlaKm!Cr54q`AOvxHWKNQY%S7@ zSwHekU@MYdxpsiW!1TnGgRq*l7!GEfe3J=_hi^}|P<@NR2a zgMZFghW~~G{!TY^Hk{F+nWa@Li&HyGmoAn;-TpGP`4Mh&cLzcyV0|wv>4=rxv6}wX zU%b+705X!RZo&zod;k)F7Y90y`0C#(*`mT~Z! zhLE|4S&EQ*L>lwI&@KoxIC63h=eDhIz|vMi)69`r>4x)(lIadod#Va_NeirGo28rh z3mau@U%U2|CLV3V^%04kP~8b{cfj{u@xGIksHFK;8+_jy|7eDP(JNcR!YivtXNQ0G zYW4Wd*gsv4gK-}!)!ZT;9!hs(Grt!$^pi@EPRi1f{c_|;e8FR-InQHQ8GmvX4I6!n zyI@66tm-bUhUmL+F0AwAh|cGNoep}pd~5$5Z|}dmc!OiuBv@WawLBdeJ6%Rb7@h8# zhM-CEHxuGU(qxIs*OEN|cOz^8!tX`&Qba95;_AkXQMy%mzN#;Wew?w9F$~)p8)+0AYGtJ*OpkQ3Ru>g?XcWMuB$s^Qy;i9{NQO>T@Y!kyZ^Wce(R2( zoTPi_jTU&j8QwL5)whgz^{>E>o$zNarkmc9bJ*Bx%5kD4T`b) z8aozAcS=Rcu!o%h>;h#|ogJ{jlVswz$4IQ}i-nzJ$cf!A$qR8bpv|!WhkZLxv;{8~ z-}-XVw1Xbv%a-R8EDr|6@%XCQh?`~1B4FE}V`M|&&KUJ3c-$ss>8>S+ScYiM8C!>J zH<8Btk0%3iJ~^kbZs(M_cGyg>Y$GTc+pt=sSBeUx-{IKBwqAF7g$8a8%0|Z*x3; zPir27IgFyIWFECX2V)@B{; z#j}3qB4rVZHsNe!$K$~cS*vhzm&+R!Q=iTne|gt%`mg25nBUHOM@*D)F7_vIHV7qP zfJYJ1FWWA`3v`02_<`PdH5S76I7J-LpOMdS=`Cl8nzR6w+nZMhIUe#w`3@WlYx_*W zmiJdInS6LgHrBa*IDL!Tomg$LvzqnY>K@q2-Wzt^aOmnL=~82F43UX5vh6Sp zyKYA)PotiJI1Y8qM8P^7@oH1L341o9Vh5^xTbz&TUK8Bot=%&@eEHi2PaZQ`vA#8D zGP7?XYa!C-A@OePyba#Bia$mzM*LdDuS4=iq;5cg8zOih#mYv_@za78pX4w7AalWU zp_3~YIEM|Dc}JVsHe_AtD5Wb=lAwe=yc)NxH&ye5(at@LV2xHM$*yviYLN#Uv5SDQ zikcn4F$j*^ayrvCmfptcuAC+7?jSQQId8U{r`C=|!mTpljm<+g1{u{X?~0&XW#Er9 z)|Pg_B4;dVkL4_8J7ZmU=~iWnl0y~D!_sx2nuyUXk&qO`_5 zbWBGP0a{yW5-TA}G;4V1s`dhr?RN%~Y$-4@AxhUlqVRTAHxV`oS9HcQesdS;V`u*i z<3^6=Q&;VB90AFgClDiM zNRRgtb_cO{eFe@2cZ4RT*mmyEspH@;UW=EWaT?R zA7>@YK}LU$Am5HdVXc!^AZk7`RwJDw=01*>QwB9=kg~`xeJ*&&uJIA=RHmXs+se&JJq-d{P)K>aZq)j(a`aHK9Q|{OGlFnzg%GnZBvbKkG zjjrt`*;Kb1I}kWsiX+9O!+tZi-^tHEYmv_PPPBKr1i9z(BnKJM}m)|@H z8~72%vGb*_WViO({#eQ*)Mum8wdElIWbAzGc1Q4RgtFzpf{UlAu|TvgeI;C<{zZUtH5eP_muKV1_;6>~o{XDU_Qy zkTh6YksMLtU{ls2WZo;ok(|1-j-wJ|5jh*%SPl=x@}4qcxZlI6U=?2m4nrgHc4W+w z6LMp2N6KtjO0Z|191+L^0Xh4zWFB5j9Fxwe%k!{*QyWfRWq(4!dW6h^?^MJv1X_r_ z?k&=n<5XCqUKkMF8L?dv*9}pe_vj*vYq%m~Cl>{{$UYvPt+9jBMkao60mxQk^~Y96 zY;nZq)*?#T*b18*u%3gM>{M?iW0xW9K;yLJLDH+n{%5W`W5q{*T*0!jhe#kN9J0wq z(7mJ^fhnA}p~_zNf#)m|jvNhS$m+v6QMbTtxO9QC>x?52oW_CD6%=9wH#I69ir(<(@?n3IOCWH$t^_QA~}PMyQ{Iw==?h8T%Hp;1bb)Tig%B! znJ8T$qblsPjhco79u7gAFu=_smP$Kx@75NqbK~bVW_OfLkhs}V4mY3ag|u!+HYBkN z61pIs3un5B#Pb|vE)Q4I9kJaJ%{r(DBD%?v-Y_mA=qN(^!R=)d@=ng`;k3gcG8W8o zl^i)CjS%UNiX%%=Vxp^Z;EDAe%Ti8XBt)h!?<_9KCI(Amvg34aFM}gn7Ie2&eXPCN z><$f|Dj_j@>DTl}_7dbQmgdPa&Mm(g3G*cKgrZ$ z402L8;}*}c2%d(}8DdNh2(IZTGxk{zve$0oK#7a^snmNo+}u`rPoof zLs2?dP!&NIIKb*@I`TQX$3r=X$V5{PtjnY(4%ttTNQN9K>_H_d zg7D%P(>U=vTAs~D2Bj3@V2R3!YE1XM1v{rBY!+)ExH>M8yD<^Gc5!WVjwZ&cIZ}y%g9q z7+`o=KcnpBaV z|C7x<4oR>#oBcerJPTWfpzK8T;KcPMEm?wbP8usd4uLJ>xHiX4*u%!IHqIvHpaU1q z+>YqEh-8&_r%2MhQkHUff_*-O&bADvVI1-4lyDEiI17EObnbC_YVa_mbKS>HawO>G zHOjWJfsc@$iR@5OZ!9NjYl9wOiv6AQpf zkG^a*Fvw&Jf}<6L5;j>zlvu=Z@(c@APB;vjj+lE8%c&jnkiFWW!o4dyLRa^}a>IEA zWq0Xi4ZKYnfSg^-Y2m@6kTf18)1)3}Z$sHs6mv5lhB18Q=E`Wb?4&mbhv&)Cb#^0g z7@9>lJ7%Rs`6F69Xr}Zikm$X|twWk_8DkMYP{$u9Q}W!oAx$G8rl z#I{E~mmYJyn;|hSh~TPlCq%RpQ4qtpmmrt(5G6M;;}R|6YX#3X(jM93giS6891eCq zb5JkrHkrl5DH9Y|b+r^Cl4Eo|<}NlgYv zj#~rUR_P6c^MPsMfXPzjv9$DH*=Y1zk|xLaXl?fGtn4WRg^Ro5fY-oh z3+6qTG$nV9Ghr?5CVV$=^5JZm>_i9QRE-^Duyv%cqr~5XMDFIn!KOLb$vKc*rO19i z22r=j`U@^bX17er43u$--L2R?9*?eVd0{@z%|*>zoMp$(7-UaC)hyJ^z{5N?mj{tg zMjlHK201fOxB{mFoRj8C7tsEVcqr6`9eNQn#UG>ZL44z0Hy6p}Ih_&J36Y&Zl*Gp^ z8k8uM3Szj6Kx9Wj0GFS0D;0XBv2b#0Ync$@)mtVgM&2%sEdG~cSeng6$V~W8MerTi zc?SY+6RfQ7I6abs^{gBhu#2@Xszb&+m^AIaxQQ$(IsZ`jcEs}T2%C=F6%O%p;5Ak_ z@*N;Gz8&iKZoFziXLRPP+P&#PPB=C-@tT(Z}523<+O3{hNI? z6hDq?jlv+1-bOT zarh{rzvaN(pNpp)UV&Y=N^6MqBmJ?IqZ?arIjZ}!=>y9)$RNnW$-Q@PltIk<6MHve zjEJQJa$ho2r!7uifNke=vd#5rwc?1|ycj`lLh=rbIt_u_o? zO*Jt$9giHbKdAp6j~-c@oU%8!%h}Sd!nv08xOd*`n~wJXWbnt?v{{@O(8aaIJR{m zyaQd5p>-9vR9Y<+Wj0s1uaW^uR+k0K8#6{EcQP6p9pq{+1S&K;NEI2OXtD#=2?0*< za*!ZLIC4TrM;T+P*xvKLlqrY9$L#j(Q|Qq4Pm zv8C=^*cfCS$YvUw*b!W}XcU;7x5i1Z?0F-{r5rWmEV!Fx2!kAbZ$&J-M!8yQvecF2 z7{yuNcforSYGTJfSH9#@%I#U}JFFim6NVZ0a%6eONa+{a$lZ9j>ejb2TR`l+58^KL zY=F$fKME&)5jXIGP0ebyG-Ina>MohglRRI#VY9d%d4&vhCCx$U8XVhzmm@pB9^3uQ zx<<`0o!bR>vWn8v*^ul6>J)Ahm&oQM%UD*LiY6;cqdVQ#@JBC>IktmaM``pulzuzy zc|2}n)vlX!xA!jC(QA)azr(vmok^Np6*i_Ob^7B)_dbw$$Kl{1=_^{YY0tgu%rZ>?HRYFAaR6bB?v@QY;_jheJ7up#>I|IVS*6ikI7*Ve2vr;LW=t1LU`XU`%-6|OP_Hx z=GGGbA^U=d9SR%$K=z%NGH0AlyzS9F3!g1pdLenrKA*lhYum6}CW_&rTi1e(?UNTZ zi{ms=c1=$gAKbuwhd4yV#WJI%{m4R9b(M#*caC#8CQG48j)@BpzZligx4f`_<&!1L zpDtN;A#p++TYt-uvBa2<%JE7zc19ZZ|lYzh#ZIluovAi_b>2>F3lXdE%uJC3=~@&4jOhYValP%adgIF-$93VjX9fi+oOeZFQiX@ zvSjfiMfX0OH|J2u$eaz%smq$PIp#eP`5RowF>^(0j!>}PjN!aTE}R)H#VHB0L?u>o zq)vriNly%B9}#D`%|RTug*pRmH7`*CKAJ6spfe zcl;)0pe1{Z<;Wb%;kz3##-=~5$pTFlXtF?)1)40-WPv6NG+Cg@0!*{ zfhG$yS)j=RO%`afK$8WUEYM_uCJQuKpveMF7HG0SlLeYA&}4xo3p81v$pTFlXtF?) z1)40-WPv6NG+Cg@0!*{fhG$yS)j=RO%`afK$8WUEYM_uCJQuKpveMF z7HG0SlLeYA&}4xo3p81v$pTFlXtF?)1)40-WPv6NG+Cg@0!*{fhG$y zS)j=RO%`afK$8WUEYM_uCJX#Ow?LfR%Gj+dc%{U-W(m)3F zsQ&AkSl2ovX6qVziFRGB#Qr1png1^-{1tvY)@?0CeUH-XyGFaO)uQsmxv#U4m~HE{ z#BW<~^VEFHr`HwriX?1XZ7-zBM2+a3wi54V zYXlnu+JD3*gN-q5o|=uO^*ZyI4Q|%;r<9@rQBf)OAY!+#*OIV(gA)5}OwmuZVxLWM zOx(84UI<7fYFPG)&8O@$)R|Hb)^(Z}6&tZJl+CRI^oD21KPk1o0jlwbYkKr6+l8x~`Zx(wGWI!w_wd zDd-JcnO{HEpH&3atulf(4i&Rm%;|ONlG)&H|J44QRUX>M{1aN&ymZR{rVyV!BeF2_ z(ix~Ecy6jw>U@;wwb>>!vnh_rJ2og&Yny9>>lB^ACWYzzxDw4vf1yrkK>6QnW&Tv+ zwha_>ioI^8Z2weQSf!h-t#g_audACGF|+n2^+9!O!=iy(g467juCMx2n&W?Ec@6Hm zN`3eecdR!=>IEg*V*jnekN>wHvCpsn45cVj+a%ujtv0V#G)gKu@kSJlOe;y=TM|8O z#h&AJ(Ofq}ipNGJDko9eeHaA<|p3V)X5Edb7xrU(OU-n#73Xv({t=8*Q zuFH`$HOi}s*Cx$HQ!^{6-hBP6j%il(r_5GTx+0}sX)`tZoJ=uixDj1{iL$)DZCKHk zzb>yj&?rFHUXy)XDRpG7^+y|^bs|ggantG4l_-vBUYnI@uWVS+ILwTw6JG;j&a4fn zY^86N#S|KycwMDFl=dI2lcVNkR`^+cTZiSZ1=K(W#qK)STEEkPl%g@&C^k=1ecctu zbYIsrpDjuh2&J1Xe5p8O zqqe_R%;0q5brsD*sBF|z{oEQ58=5NT)kKp!m@S#+?WQEvOM-LrBOBdp zG_!@x=}NcxpuV-w$rKu#8Pp>iI=b4@j-xG4#q38q0~MXetm{OLjEYXAV&<3?u$@wi z(v1YolI7>7P#RbFOkZ~bQl#p<)_#xBRC=1fd!1s2tHGH$itg*Kjj!!S#mrHsYlEwg z=&Mb-zWv`6v&$%5yNng65ql_WYbZAMH?vT>IfJ5egV&5859N>9fjcOQCd*iRF<}+A z)qpfIvr_v#`o-9m9#FT~z%^slI`b%+G!e~iW5W#JZRtL?iB@A?8$zg$D{wYn^55QS zYV|3ai$-&UM07Tyxk1Sc_M|9|IYAy2MKdR8J0&}Co0e;_isv>U>PysHq2Xo(Zm&}` zUp`g3B1_#y(FAQwm+7~iVw0$DHJv)2{@)Wb0(MYrzQ>waTW#5J;jN~l9_d>fMbTlR zQuNiFO($mU^s@TnPS31hk6PKD6)ch&WJPJu(d4|>Fx3!kGSmsN&O}YtMj_U#DYaEU zjh5xVgB+Fh%-1^2YtHar>9n~n1~nQ5#eAz$%%(Q$G`qUhioR7uY;HBCx)krsV5?;9 z^r8H>9L=Qcb^0k2&EBo@WCrc1Q*2OPB~@Dv+h>pz!x4DR4D=vJeXCNKXilk9C*?ZR z>r1qOWzS1xu%Xs49cHczkSCq%x}qQ|f;JqpcY4;!4&{&fO2^bv>$6tpwWni~LEYMC z(zqQ065QqK?=7e};@7$q9(^N_=Yk_Xhh$;3N>LhFbSrXK6 zwZd$f!DM8mX~WS5qBfB-FK_F!vrbnZ(YH6Q=-;YO(I8X?eM`kYML*S~%qY!?wYo3a znZ?Y}fN0B60j-5dzS?d7N2LL*`jPo1Q4mSf04jb*H?)!ZjM;EHzpCL3`l&X)IYk?v ziCV8ysWrYhT0L?N6f>X(^qf#XEs7&O(IlR^p(1olt(crt@YNC*FIJ-XDqVe*Io+PV zLS~YaV4ZYq5;Y*EYyG-P zPN6L!yf7b=uzrZy0FX{a>IPDm8-sHPffAY=~QeWYL5BgzLZ**v_d3G&5;Tr=7xLahWXe? z9n=ig=4>{3G<2P=QWL@6#YmknhO=>U+S8>E6{vvqw}d2<&Oos{3wcMez9k(XvspQ6z)gVyU*bn&Q_ z1^{VNRPrK%lrXWxd_=xQCGjlZSila)g!S~yr^I$4RfSpYNnbaQJTe!Qt9TWO1IBbXNn28Ek|105Nk!R&5|D`DP}G9s#?`~KKt(u}Cs{I1Qd8Jy;&JxeP>$^fd^6-I|>(F6-)LzO6wyQBThqM2%l zg^>Yuiq3EyCDK=w$=Bl3eqFa!>om0nr|>m!&5N{FB4#rx8x}+am_`3kEj1}=m0tVE zwj@*>&Fr*JN3;LMhP_O)vvD*M)a-W0glbW4NP@fL!jveEMKJ+dG*F>w;M*M)*nrf@ z*Cw%k%3{x2TV17THnLeKnj4kckih)%Yh&6Z>RYq@m98JzC=EQN6Y0#_-KZ$H;#C~2 zeo>5g6-)-ZqS!DS(FBX)!<8tG#j$}}%s`8xgRUzLAl8p;uaoiIZ zWENo>6MQ3Lrc7%3RueQU8lOqAG9p9`i;69jsKA2gAc{4;){@#vs$6Q|CBMN(=JL(_ zoyz9d{Su8v7lvBVkIaguri(`#MMK;ZzsoF4H{X^dL@16W@xf*hrm;J3L^NwnuujEH z)2y4>DSdZT5alYeT7pa|j0sjWH9oD_zu{(c-MDT~%I32+8El^FTa8aCEFN`9NKzV~ zIZ+YRUS?Hd!xTzVBP>d3VkAZB=BwgZ8Xu~qBrfE}MYGmGZ4`T~iepi9Fhx-_Usbx5 zm*InoAR&sFVm58G`V?A8KQf!nUf0>`ROEZ&Li1xo8W!fcDn@gF-==;`dl~6yT`M|K zLsUhjTxH1s+ZMRs>=u#%#>_w>S86I-7P$9Ci7RQHZ(Ln7}Y9CoSYM@k#<3p{twT-4a?V}SLP|U%eSvO-f>l&*$ zk-AQ&Vt%B5+6+$Xq^8A8P3xt}QA+58!Yeh3QkD{_l!~;g1S?Wwl_-w;QX;izp!Q0A zWHn-Yjw)qOLa13ZK9#6%l|qz6)_^E0X|wQFnMJ2)#U`gtH%EF}Hv_FtH^Zt^S0B_r zq!dlftY~Uh*Q;emf_qbete|%5hIZn2nSqhHKfANQmJ!8RYib2A76Kg^$d)N>};K={Dc+ zU)^esW^mLsyQ>~oi%helQkEK1MhC2=Cg9RzD3uwp6e<mH)9V5yWiwwJf_2^6vQwlqKBbfC2eRT1<|I~SC+yFRQv`{RvOg=HVl7v#tO;Fc z+K8E{nX|qhYM)~MT|U)6p%rVfX~Ji)Bfi**bYhBEq=-xj@g2@fI$Dsz>%p7^ z#g5kl*$GM%$I1-)>{U@5RZ5*qnW*8KX&R0hrrp!5*pubcS|`2A(ExfCqAeE_J~KP5 z+bTw@*&MN3Y1YM28(Ay%fNn75nm5PvTe|R64yf5_e3hAGlvtIMd@whOa){TQM9Pu; zl;gY8P8MhIdL%!&ijHZBf~fuRU{*|3X7v8d=mS}C6e)x=V)=AmW=v&PY-LuoNJbQ8 zUwR~^RTd9nehhP+s!k z{FEaFsg$FIX-5iE4i_XJ-JN=}B=dAx&Z)iGd`e+94OIE#(Y(0BIWbk)u@siW^kr4> z9>|HO?9Yy~=V+k3%2A^%PYY8V_ojyLO^u*=6-QP4+7Q&YHYgQGb+v)W@LNeOJ4Nun z<5*iC#g(OABehtsz7IDHk>IQl`6c|6euZj(>uL>X_CDgH*mG0_3GvwObP^;=UQg`F zs4mSS&ErKGl3tjmfbyz19?6Z>99a~L(4BLWc-1Hknz8#cuZjwvR0JhKR2thj5r$Vv zq-02xYP1^7pKD@WU8~x=s(eY3h}i)(K2eNVuKw#)i`m<-u9Ry~8f342LoE*T+hVU@ z`#lX zOpV^3MkTJ2eGbxkt-F0{hN!&QD59kY!(M1Q6R)f6*q*WzYt^Icp!$`Bigc@(&8~H# zksnQuqLigYTE*^Ny7ukjN38FX`d#{$xl3-m!MD@9-1O`Yc5g7 zsv?(INt5nbTjWM7jznn)QC`!WQ9x@8UU5=|QbRP{ncV9T+S87gZKJrhbEe_=sPR!z1P|t?9x5QgbVDT0Q4ke7 zmaD8EPZVdK+>=#Zl1)XGv5Mn&&sTr(^wAFjqt+sq zi1}3KVX>_GE;PGZMzN|1tgrjB;_4BpgX~2`7y8-&asAU;u|p=5^;4zTyPo6`5)T%{9V|$oROKa#NFB;z zHs(N1bd^bJnOdVG1+-3-8lnlh(G6AG`$UKuoV75oZ3tIp$JZkgMc515Xf~pm+9>+X zF)KFI)a^AHuB`#}AE5B*tL7*W^E1r&6vSi2Srmm&xt={oGNo8KetFsQ`TdrUFInEd zVEN1W-(Nrd;|s^Wdg{|N|GiuK$GL_<&jnnLSifWvy64>N);(%IMI735OA`lBQ$&TJ94p^L}6s4ai$vIh?d$Kh9 zL`jww-d69;(PHKJ^?jDVUa)+2-tzG|%SRV1ADp+mbM}w7Ykqz8cz)_p>3NWPnIRUL z(wL;F>P=O%f02j;OGdYYy6z`UL#*|$Q43qYO}_sc(Pe^GuFW73 zpcV!aTCY>7HM{NmU=N5Xk~t@Je|Cx%-qwDD?VF>1j(>m(W^_S%xG0mNypJ3;zSHFe zlv8E-hFHOoCJCM{&!cFRwH#|KU!JjiS!4O@8Os-EET5gRd`g;UET0hN8OwV$mUn7? z|I_K8UOx8alLtR|u;PW&#a9mIUMxvDTbz2jD6N`dFuNCwek9eyP+QW{+BBfi6gB*m zuBt?lB{JoARIXvI>Rh|4#$=6jBsT1x~XwbEIzN_{~|1$a#^k&sn}YXZiB1<*&S+ zvwT)#`GoFy*78A(bS+bBy|3R2dmkrRdss{{ zzUP{coBvu?K$Ii9(~cIUA7j_|p3GwbNDZHgfZEAnVy#OQJN3x^Q?)7bCjYM@)sSwTnH3FGDKxlA%in9H!(a^3P ztuVev{k00B@$_ngC^2l1Yx27%6*FkFu+^A{T1E;vynG<*d_@|4 z^W@%)W2Ko#OVTLEO4Cp7O*>J_u6yy>lO^dg7E>xMqoXC6N0jVIKe8wNIQ#YuIl-Q& zt4$}VO3WfJ^VO$UTUb{Ui=`v$8Y#(POh{pPudY-yP}S3Q&DHQ7A|Qyv#a218CqwIp zimY%IDgELZioA8o`;-H!dQew|Mp;@|UB3HNMG>WXUm@jGzhZ~5_p<)@2=T(JCj-pbL6)!_J@QHNS%A#%K6`GTz#Ll_RR4*lZ1<*)Zy zJ~;c!YsWu*tn$r^#m}E9eDZkSg9mfYS7x2vmwmb{g9NL~(oa`pR+nWSXRKS6Rb86F zK0}fs%Hw;pj_%DoT#`=c20QDOk5s$KXbwv-V3obuC(3dtwNIsOWfs-oV&|1vm_fD3 z)VEPriUz8QUqxBVkwl133tv2W_ly(`i|>zq0}bWsOQ@Rdx)Ri-DDj=%SA3>&4@KA= zD6Xm6qlEK6qo3s`YRBc@eR(AO6vvZZ`Y zKV|iL+48L+UthL-eEHvRU-;prlOH`?`O@ja#}DN`av=BezRdIG8KiltJga73e$9bG zf;&-`U0soTW?v5Bo+{5c!Rob)BA~1H(%`kEi~vAcfQ=^AhC0Q;@v@xL<r~0K`UQ2JCXKo#g5lwrCAhxt9(smGhRIFyQkP$36|$F!ZhKt)`JQn z3(~U(1mD?%CFc&6nq<#8!|${wlA}VY6khA{N9C2qRbdrBy(FbBX?}NJ>PnWhjEla% zUv$z`BIWJ<7DD{$vV|ReUl=t1bo!g84}N&L{MG8hCl6*kQkC^URnED6nP|`&ljS)iOO9vuWt^!@J6(}Ztb9r#Bl(Z(Fq@HpBpPved2S8Ap+a)1P+a9yPpx@s zV|5+0Q;~4ph$w8X8Y>-~W;gq_;lk4TrtsqYK)VfN`E_GtnqQJ|BErE=FK zYy+JV0~`^40Wsm@=l^q%*L59R3?LHFzs*_b z?-cFlD7JR*=G`117!Iyv>u`L22Q+!(|KzD*UyO|VN=L27-xy1I&@i0w?0`=c;sr|* zcmp3#Fge_}-Mo|GfDt}0AcLiZ`KE^Z@s<|h5Bpz^p!p`AG5wRK7%mrV#UVF^zJp6b z)`BAfldm^MyVi`?lvGjR7GwiT_5quUwIGw&s2hi^~;~(V*H6(i@d#3x_rv*4?1bd_f zdy9gQAiO#AjOPlagkV=uu(K%GF)74FY{4j_xa=S7$zqZCG*(spY2tDF_d%S1%l_Bm zP>yj1pv?ZDz)t}F2+!!x?SIegm@xcK#~i~?^Ku1>1?jKD^>~8tFQDPsr9}m#M+c@y z`}39|4EWQykMQX~Cg#XTQEWSq{Yec6JNWV^03S3hlePch_&kojN}#V3+4d#)f@3Uw z5yy7ealZ8@VnyhaRN9tGo0H#F#@Q`%pH~z z?vWhk4#Q4TxaU6>hRN)S?9T+iA0iw>)S<>Z_W9HP+;*IUoj;FvOaor$cWnR1$2&Ok z;P_Af*mlYe79T(hT_>(^Zh*_~Ae_Fdg z{t+bKB`|`p5AE(LrpZ+zV8<+;zK#<|RL}uObYlM@$WsGl@Z}NnatO`{Cq)ECyp5sH zglvW7aV&j^WSYYC?Y-2ORpE^p9>qe(+(@U)2&eRLCooNm@<@yHVh|&JL=oOF>>y>N zNBO2kK!JO8D7OwkID&ot_@9HutA&p9*=c_m|1Adh-{m&{Cr=sFfb37#>!8M;@ZazL zc~?xRM6!Ix8-vK<1Hr7AKxY4DmlGfIf3nb=_V>U2 zzU46v0Gc1|k{Rg&l<7jhRH1*eFd!wG&-lJs(LPy`o*9u|nH*DKg{4LM@Yb;p*07G_ z+&)hGm+?AKtOnV{5L+6X3~W9y$LIJ0)}QVvcCn7SfSVN^U&jvr>+Sb|J7nOH?EXxA z*|CBDSbR`ULU3+k2(tu1USjC~SiUIi&m%Yo)*e?ASbYIre*zWbYmcsc(ZT8t5J$2D z3oiYMIYOf%Y_=Oi#@U}7d1VVx#B>q99?M+vaU^{dMIXn|$3prD8Yzxz-lo!fDPL~H zwB`9%CAt&}opXh5SnL-9JJIX6N+9OBkiT#XdKNyda-{Q|qA%Yb< z{-@znK8T6q7#GmQI*^0==g6mrvpK-~Ebqm9cu2Y{l;)74C22c zXM9j@d~k>H5c5O=7?AN+kQDxZH!(+q(2lf-%=kK%zKY@UBZ^9{4C&B0JWUoz+hq*X zFdP(KnqrF@q{GCNdEy+w4vtI$LDPaD9FxKws}XsW zC3qCZc;v@;=Eit+SYE6*EH~CaC&n*Z=!FuVDU*UR{xG;aZ)wrY zwL2^|%BRD?C^-ySSZtsSSU3}3!8zD6F?$!K*;%o+tFkkMqfk^U00%&W{fSN?2Z8KyEyS2m|Gh z?VtE?gg5TfplqZe&XA9ocfiO0Fy{5`gAW#e@&%)gu*pN@Ve#ki;i3c9Khg5f$k`4t zFCm1dR+JptVTCE-9gGUoBmc3qsG{`fk~C&KJ8VnRqhO^O(PEC@)9JTl%5Og;(oaeB zOA6b56wyy2`WgGl%zzk7QJ#pSE5s^XEJkcEh4e);BjRYx zar~#!ME3yICF+p`1{5=a53_i~1n-t${rVBt=xY*vxE}YgD6%m#xFR*MI4Lk+6qu6` z2z)TmEKCe8NC@Kp#LhpHJQ0D5@#PH(g%=MUQZO?i=ugXt56p@a!{yOnD>h{3KT-A)4$W{`vm$pl)e(uuc`lQU(zUZTh3d-n)xZ3KTxf) z;yB#FQA`s+!I3v$;te$aa7MT!A5qcMg9AS1-VOrFCtSK@VK#<7h^F_2;`hEDqU22b zeLMACU1D2lbX`t(X^Nmo6jCA*lqH3fr-W9dhLvJWk_5#`!R2XTHQCXy>a3`W^zf48 zkbo@1Nf=xzV~g;!RE98t1f(>vtoxc5m$^ieQ<7f0U+=}p(yiF)+9 zgP81xIQlt}enA9|qaTqag)Hhmi)wG|v`0=p%mO7dgiu&ad%1zI8O`iT6xaX3pY16U zI)@qj7cPmmUNmisqz{GkPB?9ipsnHbE|@0Ld&wVfB(@bsHe`i2WJTS~jB3t_Y0eim z6~}<%ktxb3*&i=yp;l=)y;YG_58pejA0Iy1I9OL#pes=gqqHZQCuM^Kj^S(6uCogY(~ z9|tSXk1fp=RuziMbK=W$-xBH9B>F9tOQ~p^^Hu-< zBuxZcPM0DTK1adWfe#urk~7fcT6ti{J0N@zPKw?vI<$zZcO%(;i-5L<(v}d~5=!qx zW7!&SW4@`3epDQFzexC~B;iSU!o#wddu76goQTHU$fkl&SW8hu! zy%8DGH?j0L;&2*Eqo0$RLC``Sp-snV!X?p*LzayBr@lamA@)J&ndo+l63ZKO17IKs zc5DaDff#xp#=T676xj;LjH7MQ^j;L(!7C-LiFg(Bb6xDG`uMlC2`{UopHzlFuZd_Y z32iG0c~Tkrs66CxY3PGO!TrLZmOQ_QrQwe%q90YpKCcx$sgAu{9M+r{bSpQoF*oQ& zR$yI5U`<+Jc}hTea$sq45Ddu|RwN1l1-NH6OcazW3dl?JD@gR=EkE8jU*unu5>k{V zC`^IQ6u|ORL-SICVTGw7Fc9R!Cbnhik+AZNsEW*JM9%75;q`psje=MhAhwi?+GC@pgsic2#OaNkT(eVpDk%{H`uP3g@u%OOoqKQ)^37t`{d?FBa98rSKfjs1uU~ z#CR^bxa=>%k#Ru>$6Q+0`Sc88G@YE8$vB=>Cnu+ngvS$+lWX?*=km^aH-h z8S*J~5%MPM0`dBXqbV*)xb96yAmuqe0Zp#ZK-gqSFp4v7pO2;Yqv^v0rgQGb(0j35 zM8!6m-VrjaEirhLwLsrJx5TR z6OOx-bSe5j&I^|ha+LOw}w>MH>)FwQ?9{aQ=?r~N8gUW;l zm7=!tgnMN;liA(!*w*rRcx(*nLZ5iU<;jM+xQA!I3RZBIO|o!uz6G zAO$<#9>51dFoNC<`^WBt(0hTjEr>o0rEOs>S0diFgt4I72FC=-D9#_F>8&W*1dIH2 zGxT>;C@#JP@fqBj#U{U6|nSBs9$Jf$Y2;~GNxSEEJ@=|NDHb=^Q%tt zsY~;zP4&K>=6^js=z3P@_3Ut{m+IW8n!LzbSZ-)-Zb*HupdmlvMt)R7Q4GRnQ+eXu z>V$jM;28hndeXbPluvaDcy<3tfk#z@$fIi zM-m+Ojv~89k|RoF?`UGDCHFbqlxY|w&9`yP89(AZ5WM~Y!GkiNlc`vL+&}X6aV?6zmOi}Eh^Uxebn}X=wa1JGxB3q+qOBD0U+tIWc zJE8PuAiWtx??tj!65csMz<-^dd3n3&!nVihcHgT5i`{k8t#>_1wx|!%`BLJ+-1%4y(!%z{RRE!YY9hj+PS48=MXX zSFOnmuFLRm$nd(I?eVz4`)!&3$Ex5b1)leFy;`z;8ngUrv!R^BYap1j!{CPX`QZ)u zVT}dhjRld7MNv(qaZMF*O;xdN4M~sc6JFHDe5i|fTkf4|zvlzVTl@ZOVh zbzS@=tyrrKnXU)gQZ9ci_5AyK;HNUbnlRJWROeR};{NLCnJjW%SlZ>vs-A9}2VP#@ z6?S!XPadVsn%?-xbb++>nm*^|N+3}lRTiT>s7MZuBia*5p4AH4;K|yvC@@iGzQ>z9 zsPs{2pnh!kL_oX`77-NSU~TXq0lf#O6EK7KVKQJ2pp5;NM1Q5SlJ;Br?-HBag5$K> zgVuU8Ol{CPQS>fcID+j!4k5{+w}2J5chFs2%7zB}k^Up!_S*IdSGTEzoLN=otDAl8 zNX+@QUkYtMzSW1aM&f5)DxeIGy=)Iz# zn|Z;Fd7%IYH*6{hzf~C4Tom416nU#Ss<|xoR%Lw4^@MwM@ei*@y{;2HEpVt0KX!F% z-}B3e(QFcMaGM*p%3=EIb-st+WnM|Sv@ym)tJq8Y|Bs9a zd1@vxU)UMuy=#QMR$n`f-qtI-^Bf_55OU2Ib%JA}x&%=D8RU>M*)fvr8buC_CkK_- z)?!%?p~1|z$tLu16jQBmMqp(eqnJ{C5Wx)8!0v7cZ4O}US_A1VZ`$HZ?+4Sn{`9st zeGx&wqI=^{Uxm?ozO?w#oe=FLv!x$%tiKgIGL)DGER%+@>wiwAFQZvJ|Cz?Vg=hfL z^%X!-<3zTGov__L-g>)=`6k6Hd!{<GkM4mBL%)!s|s*)kV=2`BCMFkh#cfL3L?fOt0@oBvNn`EQU*~Vx!Jde?)5B%w)K-%O&uiMfF=U;i2&*BUUuCA?e-}5BV zxHU}oxyYD4Lg5*~iXGs47frt<&>v_P`O^Z2E> z0r%1_N1DtxT`6I*keHzooI&)ULxzzPisa;Y;<2fh$Y7H94ta<5{Tw#RUDoUlv(@T# zZG(*CroKGK!xP9+6=E=rST5}jlvc~SUtHD|u@nsf=vWk$@JBv-qO>$$+*xlr(VIz5 zK)X&Q$5fd=B4ZvH%dUVw!Suv5Qf#k98$IbG1RQu`7%TR%Whyk3g;3bD$hRq{Z~E`} zo@4qr^3bD*gYObg((+5R#Ddn^((6}gjV&#*e4S|UG*<71_pYpq%gP-$R61{Z5_OUV z?XQ4{u3#kniWQk)`pk!Z$JnAv;6>{lXpQ~DSUr!!(=KdQKEH9o3de4R4b)p~9+w`Zu6X@TF zO9Fb+@$W)2T5C_+{OE%amUKR)(_105Q2?FxBFE`|ym_A8#`tqHFRnyXB$pNwE0hmY ziN#EY^5i&Ts6wn36R&N3y|(su*(7VXzQ^UYov&(jvDNBqwH$zWj07CDQ5&ifeKle{ zM~ZP=+Rb#nBxDbQs`cs~*s)pD^V&LDdo3B0`I7KQ?aAcO1cnlBh_6HQ=u9HktL1d3 zu&M+{_KYQNyG9_pF@Hoy4)|{S(*}213P28+$)B@KYh2cq+pqkPV(@pa2`#q#g>4}q z+B1Tc*0j{>cb@56k?!Yoqx*sb55o?%`0uH4THoZo>wVIhFRABJjn#9lmay+W@}cj; z=?946K$>syE!F6I_PIwfCxdlopI{JtH7f3)(mS@iKn!kg>yud1T)?;SqC2b;s_O@G?# zNALIm4XwFEpZd~w0{Sza{zL=HgWh+et&a4e2mK7P(e!65{T59duFxBo>FsOumK}ZU zMc)E3<~9_aKfUSsCFkP3aGeCBS>8K_IIQhtuR$&@ASN@2t}@XXBZjD_M0O1%#?#5A zg~VZ1X9tZgE}OgCXb`KVAh>1t)>iAE6&tbWiG5Ug&3k&pe}K7X%AZsIXlgMm3s z&NZe5CbZUu7Ms(`OEkyilgRK{n08yxp(hb1?gk$%y1Kr=W(_T|qOXyB{pcd6xi1q9zGa>Lka6a7j(Mu(3j3`J zl^%y^EmAgGyD9W;EWH^)Z-g;T+XxH)QR?`k*oM|P(q=FE0zFEB`3y_<70`@Bw{#`=p%piCp+Yb*$iShfo(4>U`(y%kt>Uc z4g7bmm?MM2aRDH-SV4*_C#As=24@6QbahaDpmzsRbbK=72#7HJiQEeRg*w7%ILC6y zUj}Lf2ELH8=1aSONHar>z3I1?me~9*y!bj^_j~49T72Pil0G1QPBnZW*w^a6Bhz$_ z$JT*}30GEiciGxM)$#`V zZ|~cbvtKezUPSBO^E>c8)wnHWXS3ht8{RvU%~#rP8Gm_$vYEzMyB$-!wACXF=YGvH zeH6C;nega?u>HkXSGn$xzp_>qnjy?!cE0_lTOlXzMjC!Cw51g;wBCnSI?_rTT7QkU zc{5n|J!!KeZE>Wx9heV(0yVr96UGuWwt$Gpfc*HvpMHy>jn=gOGHr9A@5AVuV483A zMRb_fSknAccYN1|9qw;}5nDj)R+DS1iQP(aW-2j8h#N=rCy-;K$+@Wv7XDn{gX|hW z_6#D2hLU3=894i;tiGz*?P>I3+UNxP zk$(DN@ZKiRt@#(0WSK82wp)AC|3I0`mYl2Wv#d2+1={uQ+i$q+U=AC&`K568o5bU- zLHkn8R>YgE^wXYodDG}1J$3i}q|Ul-60XuvvWeQP4)3b%Nfs zrbVY|s{?)ROJ4;h>yAv+8JeOuEa~L1VoTLLQ{_U7$#FXJc5BFm`Q(HWb3luE?YJDZ%sjZGzn(}=-zW(dwH=VpIPKS!_Ivl*Ua`Y>p>*VcZfaCUH?$$Eo` z(x(tTrmZ+oTPnxacfBacn+**!c|dqSG|!ZeG;sTY=o z8%&Qinvrx?J;7l5{lMMt;}4{r9anT|W{lqWY>S1_2GgL@b!QN(wK4{?$dU0xTag?a zPYkCJy-7qLvk{7Y8aa*ooJLNnkP{QxN5(UV=}cm~l6i{z=FTxE zsWJIxlS23Rb!r5*&|PO4`sL! zQE^$Ek8fRS*EzdJlYI!th?5)|1VivbkOV|5S78TX7O@Uk5p@9?G8Hlg=?`Ti#vFnc z-2>kpLxOe>_uMkTd7TVoGkWY8nRDs{O$G#N^x-k>fvLh?limsVk$FDqxU$t^iL_LK?59DA!pQx5xzKsoSVi3#wk?>6o%u5s^lEx`UGM$nZ-cZc{OqwrP*wDGR_1P zoP&9SpSU=W*ezkQ^E?&=r?4x41#S!y<=hNnF_)NQYH=k7cv;RNr!nlR?1(u4sk1Hq zM6SeT_YY$bL0?aW`JmHEVy4ER!{D3zf4c#C41v9qJC6YKWbBEFA$v#q?-+_PM3lI&K*~^!n9Y|$j|fQ?S3vm1 zYzN1Uaq3T%P|A)!(iWWXCYFSAjvq)J1#TX}uic zBs>LN;4;RD!?ReN#Mtr+So_DaxCdu}AK^?V&7N-SWde2%xw^dbg@qCp^NHzfVvL7H z3RpzocP4WgH5>;xR701Sv1*qkdUWaWq(F@E#FU=GNP!4k2}6fjjhT!IiR^$f3R(^= z4(vb)VD^HM9`Y}cDKbOGfs26%jsw<_QRLWI@f9$ufe9R8*umi7>)6*tlEO`5ZXlSC zX9fqvkHYZlfC!X0^9*+nh>hlmPQZ$IIFC7tu`02dPb_B;Gq^2okMnCh0=x#d0glLV z&?;v!Em2CK#8zR~L@-2P1YdB(WWZd)OaocVdF)Qm0l+ijI2;)U$qL~Lor7E5Ha2~zLFOZ3}o?5*=l78nb9Fan|Y%D6lI_i4o4k8Z_=-mC-@!K_o$Hz^6DFM?e`51t-GH#}@lI zlZ^vL46~OTWDFn-55;dWt52&kLrH~UVX}M@LySRye;pXXw%`ay!1j;=cYPHx7akcu z69DDtDDh=+1YCF@JPKDG$Ho~1aaj<={lMU57+xH-S|H|#VFX9Sy>qI>8uJ{3GnsKQ zozAAywtmd7O|ev@&ITM1al@=FGlra+z}(FeGkgjcZ!vpS+1)ZP#x>`&%NkE37cj$- zA2GEsKk+RUvbhi0*o$oJ%`Sdafpc>Ny}86@331aBbA*c>!{cLeBGgPEy4b;_MI3`~ zPGDojO9I3M9n|=kREUtAA=l_n7RNNm73(F$b~!Oag5aD8bg*+=iB0vxW4N$@pojE| z(N<;wA3nlz6Ik%BF@bps9PNk#^A~)IXXJ}1P(y41IBY=#sUN3d#DD{T;tDu<_aL^# z7x5GDP@V{GfXBg42~2f-5nK?XKngd;Bg5oDy2KYToxl|P*ufWZ0*2Ro76ZhOGtfZ) zAX-8$g5wz#VuRp`qHiJ#iKrBChxje9fc!uP#1$|-4h|#wdW(q%Nj99 zL)3$qEfd(IVT5@@gh9ySV-&Isb};7n2wd1>!XGi_7-vpPU=H9o!WsfNzIBAlLHHB< zI1OoS6dFs#z4(?cXs|Xs%Up&V>g9!{L01%EG9?5Db)}BK8MvfT6 z9T`vd4q>O^i|{D`#h<91afcW>=IKbLNR+q&E{?1>g}DZX7uV#2B##LvV3hFSfXE^P z=bre}5^L)mM{F^?&@=E;9K-;c%w)svqDjtAXR6X}F*%ECa)$W^;Ssd0GvXoS2U^^4 zc8uf!d6vsgjwC@FNXeNIMW^M1Hxb(fOqU`!Yma0BHgIbv6ozm;1RQuA9tSfMPlQQ} z+dep)Z0N;`4shJkm(4~L9ak5VD+`F@@^+4hROlN36XGPgqfmETKLI0v#{>clKJOP8 zRRjSH8nO?+j}tJN@%a2{`-iiA(Bw(s9`HrPJzO7H5$kvm_!VHn&0vJ_MO*|scwV5y zJ?sXcq0B3B5&Q{`_!f5XyV%DyF|RN;Fk28v;mer+JY0~nS;~eEUIs4M!f<_-B@y1$ z*gb$FTo}_HXfR^DffckDEo25Gkut7t50@fY<*^zB>Wba@mlC8mU0dZJNuFey4#OPrt zFiSCR{D@(!FiY`tKzUe!4Fx{MK{ke?$$oiuS$F|X#-JhlV*)crQxJ~{o)^*!uy9$7 z8-^Mq&-*660vwnFFpMXz4*`Qw!o$MNK^Dg`58_Sx?$ayTZj zERRr$U%{uCvA$cnhV1Eup)=zeT(~#>%9tm(2Yxz%rY>d+XBaL%8v-}NCJZi#oC1!} zoY*3dAkATXkRS0AxI>(PxC&|Ev0mI%L1zrfPG>YZfqPRT;k$YWwsy9iPe25NfRu_~ z7ao*J(HUHPR-r(DP@U<>I`feYR%3*FB_TVnEhF}eh|^LMyoJQe!7jxtpVZvkF7(vLqG3>j{U*H{F z-a$@*gYd((r4QN11vH!w)|ke4z(XU+f#Iz40IW_Ll6Y`j#0c>XHMJ@^pul+HWK0P} zTIR4^)I|`00Wqc=Ukl)TTwEWY!sVeJ;YX?jPlrn#S0K7bnq1rfU);9C3gWe~3$hhZ zBIsRR#-su`+Ad`o4IBX!egsd)&EV#du#h02Ixs{)iF2$MktLrsm~$thLb30R{T z!EIyu;ZF<#iz*Ngcm!yYnZ#~63E18lG%rGI%wboe8PPK+eyCJ>jUBxa{%+-^z1dXj#q6Au#XFy(*=UW_D*&K~X@ zGaok(n!Nj9&~Oe$1Jd`vP;zu61B7!BvM`OoajP7YQb^61ftMGtPjT{Tl+~CF6Uc6P zmLFjd7zky2n8k6>r&Ijtu4%UKFTZ4Rb56O#wB3eybt38u(m(2%$YYCsePK#3hdHH-U3Vr*ak9&5;{?e!+aU%!@+m=W0NgvC(qJL~!cFzP@3c?;Gny|BdFU{jrmd!#npCvxi@z#OT*CV!Mi_>H z1q%#2#s^M>9XuLh2rwZjAcOEa1OtU>i95$g;ww0lcVrAN9u{-p@&a+?3x$bs!1>_9 zM`@gcC=FjesmQ__d>@ew?gJMF2;@qTJur;@9DWO~jU$5HJ#ZViCD4yMCWlZ4;;5rB zVX%XxH>73^J}*R%DUj{5|LJxj2h>X_1IM-2^*2`)!g$NJIb!eV|0z zG8|7L_w-CS(8psLxiXzZ?jV9~47&45a!QeyBq)Mip~y6^4FqtdiH2X12b1JLek zS?z!-2GNG5@6FWh8@o#aIkdz;=DP6!u%ponY1|T;w2Y>&dFe3Xq1_M~H;-nlq*1fq zc_{uJJcZ|o;lj`%9ba0=raguoS_}Kg8t@bZ4Y&pj?FdXx-1)T?#2kDPqsFpcF47qm z#VpW&Mn`B6Yc@g&BMf4NWoKVDO1J`i0b5W5JDdhdfZz{}hv|*c14IlO<_5fG&mcA# zAf+&=5D)#f{Dr^;Qs8K_gc*?G_V_4)TSeRhAQ+qlSKdE}I4@_K!Ap~PY7p0z#AY5- z_|SFNmLl#kxR+*;Ylw%qXEoxmgm|oBC-Ao&u$f%snr!H^ox-{Ux+`b`Kqf(qO(cS? zlJSRnq@9q9*x5O7izF70V0Z**)II&k!6B?FL>Rm{g~acceBw6irnNFWD{7bY6*NF5 zkwU!zEjG$6SI0eZo0@jKtLrLuaiktdodoS3m{+-mecN27)|d}%xib2W^{9I`3U{pJ zAKNO>(AhL?4UL>lqtt24yq|%SKl`eD@|pO|VdO1ynR}N9@*DvXBLk}9{!s;V>B>_ zffCg*T;T9fV!xOKZtes}K!XV@jlq59POEU zOeRIU53Km?@LYpIiHCZK4)+x9?i8^@GI%42+$9mSuan<8;!T zHv9-%{5f#e-=XtfxhlQ)P@z%t-#L$^f@#1=%v+byurJ;dKYLAhVn3XQ z&*C|v4}&-c1%gBmUcAOe*aX2#=&y1Eg$RJ=L+Flf#pjhXPASA`4?vDUc4oeyLTaq0 zybIg-#CyeU$9au5Gg40uLb8S|L0F4D)Fhy*PC$@xLJll4PxgZ) zY4=G4HSFsR$k88|tv8_P^pHG*0mWwq7oQrKt1pLbq0!)iQ-fegI+f=Ya6-vhc^oe~ zJt$vaw%l}RwZ({x6FrlUcF8>1EAil8(Ypxf!_tq-R+)_{JEsu7jW{kMg+>Ew%}3#y zar-(4AO){y-i;Ow9QN2Sa&!n09qN{_Pg1l`qRwI%obI8^l$!G+%gzi>KF~7)H@}tC znUAW!FbWVcqacI?+qc9c#NjeGilmVn!SdK7Jc!XNMpqXB-%R&P>5!5LqG&d#4$($!R}rWdwU@l zBL$;P;kAxbT~NL6u{2~y7ki|hh0I6bK>?ej8f~V5NT%+f^b-S8kI6;u`U}$)iYw)q zT&d}(3iGkm7USwKO}cSewaR>a(V3C?r-l`t9+9m#Fh_q-p5YK|K>?Pk(>G10KQ5J{ zKP366Z0a%DwBvo#j>`g2x=!EZBfWqr>tw$iZkwUgH}ixnj^H}!I&xXM{o!kwC**SU z`eo?IWb4Z18}?5=*d<#>Hs?g&e7*ifhJy+X1{R+g3at1j{Wz}HCr?kNz@TsGnStf! z2Ev0<4|U2r*{9w@q12#X;YpbrW}|9NMwA*2$~f8sR%1G%>~#NPBe|r5QsG;OH`XOr z5NEsrwTc92k;H@DAGyzZ8?fM0@S?|VYBw({mKn=uoah(3ue<+Rc7neK3EN6Ai{cM- z!{qeVB0lRop=pqNQtq+K)bF87J_XGF{z+?1A%>j`YS(mj1w$4*$I#YZuF|H7EyQqem*B_8O3ZNv^KxN7XaE^uCK zF>aANBohvFOW4~bs@(%XvyS)9(dl*FSib7?pa!$y zP39wSnT@E>ldaN|Ej`gISDb?4<8%?3AL81c7Y$={(X zJ_IcL9--OhIQr?@0s}^xs&GHcfRPQyQIC=qPmnuBSe%M^}BC3hGws$rCVtFG6pefCKWB9IU3B0H`0vNKf`B!4W9fXRPBS0 z(kCCK_nwNsL)CbWD03i6QCTD3A$Y>|P_jXo!KaWl5Z@pP+*gx0j4)=q(Xix`{gL&M z`gh5ay?sgCepzrVFp`JRN<7>xRl5)HB_GC6caPrPS-7V&_$D3g5xG+e89!=QCm7YV)k_p-6n+xjM~#VY`b*m7HJ<1*0u`V$et12&Fd5Yjgo#FBto`z3Ek0E zuvMCUx{e6Y2HD#!bgN|K_D*0Y*dh_OO$tI0Y7oBynSqw3U^A*_>C9uYiF+7LW@e-)kRS)$Xu>HN@Uvmy5^hc}xJsL}6HcUq>=bYPRk zh~|r<5CA_1&!*9fsBrPG@CDz3=D+orRifWN^t zNs@K?R$d$%xRoV$pG|+ctdT?$0~!Oxizl9OU$46^volZhX|-0Q@hhot!OxJXG;}&k zi;=UaXeo_fNMjei@t*j~Q~B>OC|KMBjxXXpRIGE9SdO!q%L-5E1T+@x<`WRaAY&N4 zHHg~^l(AAdhAhP=9PWYY8*ZzM0uWtLd8XmftPmaT9le)n>4+VjF!D(ExB`Zn-@=cu zAGApV7sL@5+}s%ga=k>*md=PR5O0Vxft#fT+q?K}kak+mY5~X?bT1+A5!W(p4 z!SoDv@S`vYWBeJkNy>knMBsV}2on4RM!#_#4?)O6T+l2UI`>E5v=6>h zen&2O=`kG49jXeDquKY-XkLN=4| z?b7~Q1idND1Y}8{_3-81p*uSTZIuE||4kTYDflBEFmR(}ICHMf*awB+&5{rg4tOV1 zo#_Ww%5iI{khqXrID)?7wIxjY+Ad_G2rXi}g{;kJhhk?T`$ezS-0xx%t|!6kN%T&V zd8~8JiLRxFeKL+n+`KUSuI-q6)}tRgjC)`==DEuR*lQ2f558(NY7tG-r0HvE4#Jlv zjhgpUfS9(NM$Y@-HtAR3><3pzWNVYmBY#1bAeiSF4J|Ytm7+5wcv~;G)sisGOcef3 z(9&y&la{poGUAEiaRK^fk`J86=jir8gz?ZMUhBKKX?1ec5aY)A?&~^X^#l_V;}Ex} zXSKKG=7yOlCs9|&ui?#1yf%RMAFv{?L`<+=JdT*+{Yj|b*~A{;<`RbmbK(-3y_ObjrX|~H#ZFqilV+}<@$+fX z23ot97H*>HnlyhCOAyISHZBoKhkCCPiqOAnIAg+Mx;DTiaLyJxh%042K zez1GWzAj07q|*=nm3_2(?y;T)Cwt``>t1FcTX9OR{FDr2#(mq-Z#|S>yN~(cHU5`i z>aXBQG;uCXTTD|H(s*^6y@KYfqCwXdGM?)t?NK73##})d z$K}LvB?E*VTg3m>#7R>EWjIzb(vSW1(raSVr9pmNZgPUK(L_+~1c**+BplaB0HPiD z<|@d7ubZY6imzPd+o=xK4fUpz_*f z%3Jp---G9U^qTo2aQBy{p4chBbW(leI_R-x0Y&amGtvd~9+o#2>|2@k=kFUr-u^#rsdgvR+(eGWy{PLefLnhPU$uvZj zhN#dmRT`*7BPUbA1R6Dk?PM&W$qQ)eB3iwTDRy*4B4_^&n)2OG^|wI%TfodW9#b(< z6L*uC9mHn^aY5a>oTMM=TWg_mhhBWztY zyFkDoabw*oS(`lfo$%OWB9ad4EvzP<>(~+rPH;x*#C71g%UO!VG7UH)z#_Ur`PeOk zD3X5~ppk2cuEpSLQ-zNK^J(r5gbUm-2!8dM@Yzf07oIm%jYgtRF6Q{fXZjl#m5-iN ze+0~defFOA-hJ|qfY~3tro6)x^_%|LNA0uc)VEHP-Z-jY3+}I-C*mt_T_-=YS9)kY zh8yi+cP@;)XEmzDQsJ)Em^&9o-?kVD0bg>mPsX0kIS0C>@05lW9PVDGBU5iUsNQfu zy?+1nZKV1{uUlsa)av%C*YEYfa@ce0F%3pC#fK$o4Z78z>h{`U%uoL*G-?j(N(D}) zkZPf7zkHSd_Ew~Dy>K;Vg2}(U$FZ3iG?99br9l&@a0Zt?me7igv}^;d-^B{&T;!D1 zG8yy(MdaJ3GrW2En41fk>3n)-8 zD<$k#N@Ca%2Eh^Ub)adFSPC};MCa8Kt{R=NbzVW77Lcstot}D+d*GoMzKaE0EXpu; z0KHb4J55SVA2e4n=e6Ok#>9^%n;}?*$vQpyWr$Xm;ZSstiVWnQIH}OYH8gh%ja@{; zXCtou3{qv!p0b!0t>yg@akb^bum+O>HK+Pk8_2-QPV{Ip9dy@nXu}!Vo9FtrU6gOL z9NJ&>9I~#Cwr9YbS={EQgXc8^;3Pzb$da#-#}Sy zKCJfKpv(iELePIk?pw;X0n16qY7(JIijH-!)bCS&u3y_Fg@$u-$y-T@w$xi^B`Sa_ zTk*?p%1hf3-@KG*q&gMOr7?48#I&D2<9_>%qoJxa0-9Wv#?5At6I-D=OZ~`qfvP_| z6lwSjR%#&+LfYl6r-{pG^a2{WfTpbd;5+TPv+6yY3Ax8)yrInT`l2$innL_EN#=>c zH5QW^uFR;hoSbxQh|fBvHeqNcv+`jLTRX-+oCrGzo-h=Kj4v25bt;7|PJ;%kFzWT( zRq=(7O5|?V=LC0qoU;~gC;^TxYnUMp!RYMT z6TW%dd?*Zaqd>b;+CCDymFWcUCBz2`eR99$uw#j4N<0v>MTyCEuuNdeJg14LYlsk7ObbDB{XI} zja@+Fm#|421y!&BdEtfA#0U0D_idGLUQj^T^u$s--iew(?B}Ap@>k}Gp;hLSu3u6s zGFDDHKFohp7vwk$F9sQ7iyfdr76S-|Zxse#f`vrR2d^#_BX&VS1Yx5=P8q7)gA7TL(69zI#gBdCj7-k3-2r4hCz6sVWJSAUj zFr-AcU&Fa!UwmfMtj)A&J55;1R9p%kRvcF|_mg;TY(rO*;AJFIgCuSt1xKXojr!cR z9P;pz{B2{IJLmh~w;X)?jLa<~nY-r)J+PF&Yc`P)bTEDgO=3Y31mot?_<2m66l&76g*0y!Em+O!CWOk8jWkb_7H*)~nlxL3rmdiH z^J(0|Um-K!d8#~hQo3)ac*kZOR)5lubjJ#;1>Ru9JB|3qeRI(I?wNLSD0+;I_Uh%9 zlj05!Kvu%5`&bSG8s>BWx15WjDu8n$ebvOdu>#GV+E)Kb4EceKA&;yHp4=v?h z+YbBaI`+NOIDoD>(KT-$$=ydvwaG2xUauTSA-g@iGW4~J(ieZV?*Y?k*jy@F!78Av zwcPlxp-GErm>OmjLk5wVyNW4HNKL3pU@F1H+sJ$^btz4Nl2oTTt7$H0Y?uFaFjyi1VfccU+dd_Fo&dy<6<|?n%45)|~49#C9ANK>Tc>gUW>XG;j)&7Ev=lIgfr|A^Y^=fY!4;TF&%*Y}xOjMW4r3GBjWU z6-@o+KK{;m*%lMohZhF}_ZN@xKYS*9@fzP^+UvW|Bo^xu7SnWye@&XPhGrpOYq4ce zKt$6%87riS4@;R8$yiPc*3z0Ctg_2R!HdYbmgcWVio~KDBnpk0%ZfqK@@Ed?D)f8S zoEdQY!srIGk@?5^2vL(RB$jxupf@=-keE*<_Vdxx>z-|>P+>9Yrv1D|ySXLjvEn9! zfX2VHz`N2_1j+V)xrPu5f}n#@1zZ+dS2BOJf#C=^a>EXlf({ItOKe?cDGA?7UU@0D z+4Mu#I_ofT;?~sgE6f==7XvO=27C{! z2D4fM?bL{{KHS>ljr!)zRENL^I_KSkSX8%CNRMa>B19CSU|(kNQAsu z3@@X}DBBk^^pWac{idKf_QY<~qpKsal!J@E^Pcj?bMhkxCDf0JyNC<^0TGsZ@yd-n zIns|D?oTdFBZ2EA(~b{86;^dot-*Fq?bSJj=O;$&m&L#`?^`S$G!VhiF@XV>(O`p#4}G`8vgO)xqA0C8W!vLr4Jg#*_^ z%;Z%l%0`A}jRcGjsys?IH{k5;(*|J{tZ^m`Z68)tYDuF-fTHz7)Yew#g=mps*%U7f zpnk729fa)P@qW%vFGY$LL7*B9Rc}2fld^%NZ6?=GkXN>Yq2WKf4g2gi^tZp_cP|B2 zo#BlF)L+6mR5Xt@g;1bjM>zMrn-bzDDwOvg6J9!xZn+?bE%v|qqevO~18$4N7oy5? zp8z%83@-OAV8txF3=Ls;eau1{IroFR@;jKP3c8Ce7lziG47g<}|G-x9z4!EIE~+;# zj4aaY6Rbrpp;14E9F!vm`jRt>nmBg#~Y5|HJ*-<1Ur_e5KD39QYKklzzcqdg8&5M0w$IZJeHE8lbt{Kj=62s z{b%I#oRiXc&49hg!KzCsexA6llVrA0+5s@PX%5X;L}SoqpGK3pUjKF9mhSprik^RTI0dj&kd?NJ)r#5fHK4WP~8c8q_EzHl&3v}92rQC4kl-m zu$&qLJY=K9$xg^cP!nrTv}!rT zP&n)$vW_1h#BwYvva_wscIRT_)eA6a9{Z_*D>CDd!5-IQ*g?Q#F&fVa-0-Z&|s z>A{vr5E8?t0moP0@vKryT=X+!+6OOmeiYiS40>`+;j`aVDqQp}a5@a45ht*(#LW7Q z)=HEbO21!$Dl~WsGA3RPV%3CT<~MJZFW!^B`Aq%nIpvG*v|Hx~-nlURjf?6#_Zc`7n7XffpeC`zYh6Q$?r;`1clRM@@joE53TtOQMA`!i z&MFoekGo+r^MUvB8@97j^@sD+@UkVwfp;s6B!&e8f}IY0Ajlh1AxI%t;(Vt3AxPkW zIEVm?PFU`>LQ@ydB{r5>SCD%Hk+xi4q27oH11 z_{x7dvfj8R9FWn$r`Q4_w)~{RPhjVNI5HUwcI_xxiHa$tLC~B}&>lcbPSsvk!nLirS@LYQUM79#snmahG0X43*iPxBq2a}MidYz4=aX3Qb<58DuPh3 zyu=5B5uaBOF7@;`^G{B{h3)FMX3d&2-<+Af_kZ@xo|!$TsTtkUOo_KgDx5=AKM}Jl zD|$y()ZUy9czK;F8E|xd&+U&#tsHR!+K1f&1PB2H)t`|GlJRjAlsr<5G6ZYVI0QN6 zL=#Cf4+5)o$agCRaAfAPc8EJ4WGuyM<%G1jpJ9mjYH9NM1yO&QY`o2B2c|VEcPC|b zM8yhn)~tozX{GQIh(3;FDo{o^LL0%7;f4s(!kU@b!`;a6C1$)v1}SFDz{Lce{tohVyM#mZf}A6lFnV3ur@{w~Af9 zJH2ha?E-YZBa&<-L6VqF;csw^Zjvb-1bK}|D!A?F? zkodR!xcBqh?VHqc^Mj4cG3mLI9BU`LF|Dni(imP3Q5 zcuDn?CbmPXQ%1@NK3*Z=cQcVED8W$y1nnXeSotEseL?gYb}+nmbNA+U!0gfQ3g-pu zx??eZw2II+W9f7PnC1(e>eYY&hTiflc z#aNYzS1NA#ups=zjL^EG4(vrAdp7L&%=oiK-Hy*nc_TY={cnQi;=IkZL~a>`nlUZS zs20}Q>64Yy1X(m97k4AgI~>Mm@&v7W|4+^ zwKk#^OEQv@bWm86Zv!JITG`iK*mSy#9s#S~W$I*oht0|B^9R|@ZH~g36RXA>EWM8l zTGUQ3>?BjgDO;k zCIc*5$pCx2q_1~$W;JKP?q~NWU*q6A%^EYnM<4-wtn4=TV#Vie0Zq1CSm<$l$?5=K z%I+6SJ6b4T=@0s63fg?RBpDUMn_+TFLT@lmH&XeYK|5)^7al>bs6 ze~cq^+FsF`#OrwV1q{(!F0JVDRar92{4YzAudYtR&zXgux3SO1{~LW~PQ>Xs zQ6Cn@f3f8DbEUT(&QGY#j;y%9@pD|T<@!K`85?3Ia>1vi8OO*OVG7vf9C^d*Psjgd zT2fV3#LmYfYO~|dEggKme8`cxcWjv$#&Svbr?3(MMIBWr*c{koECL#OwUZlI5X8Q) zkU^Vto~1?6qz~$U!f|kIH(N}F;)Sngv4E@w7+?4%JrP_ot@o+-D7WSm#>yY zuVQ>-ICTRlaz@~`d+^rKMx-0$@2pcdg!(uFs{%P9Nxe=nWX8rMb8(-T_{Pqyz16(mv#NLDB;th zq%X=kpDKv?Xny?Zxv{6_#MYH|IXNeu<@wt=5w(*;@LehDWBonG8uA!I^1IsPl4w&>mtCk{3x{Y`aN*rEK?x2Gr7Oo`$;=lSj4ceAWjj20zFxsnHQ-+A|kDSm8pF2vO(h~M^OXq=`aQ=L!xa3 zD4Ftb%e>p1o*Bu`@R?_SRyo3W>k|L|QVMP!B{=CpLw|dmVb75rh$0W-H9o6*_*0>9 zyar7kdDjP#G(-k+p_*Zq+5iq52>{AIt+r}y+6CeG?b0q^lqT0bpM*=;*~0jDrrvZg zC-_)?_`XTaswXt5&1_ba(R^>_jYsohKPX7S1Lw%>_?=ldaV>&-mygGq^iVS@#Eflg z$4{^u%O%Q4GmEqJ;n&qpiGA<6PCF)s;Hrw(_wjk%KUvhTZpGlEbJNz2Y>wHF<03N= zhKFb-_Jhz=e!Y|&f~5H?5EpV_pHURYc*h zyFf9ORAHY}bzmV#CS_opm^7FY`z7_46qyY%tALcGOE_bORRhx^Rtm4ou zwr+Gyf+HKJAmt`Ci=&Sk3Dhd^k$RvC z79bM7sKMt>Wb?02wl)ET8oNwBltNbexG&@Scb{p&ZsEH*)(MSsbC&0CJsr&jlRa7C zTSnaocDdaC=At?`kweXdwk9LOj1M*AnEpb{)C3$In^b2;zcW4c&72OeJ=T`XHH!wD z{dt|vE&tVpHTNB!cgMQCh+2K`>g{&NA-7lZw zih3LKh%;9nDpp>A`gs_g>jD`gn2Q z%B&b9#^*=ax14^bndOv@>sRW^mS8j5&NR>D@W_eMEF+zuET%`E5zbn=H>g*zFWOld+td$ui7~=&f9+L{wJF( zWc6{k_1v(Pu)dJnfE!dTj;=yA1VFMdHLm)oU89V`5%L5a3F#v|svnZzFBq?{)(x#P zZ-}o6!tuS|Cmha8tjP-H)*fr6soXG)v3ked4a`o?7t+~jX=6qPTO6NBF{|!v{O63Q z!+BlzJsH1kOfXw#bj#}OxW5(m!o}-EX}=BWA5Ms0!ho627hkzzHo$&5fZ+=q zE&Y8Rf=<*>Yh~+DQySI?yzYZ26-KE6i%cIE;p;Q5SBQCCRUsl3CYG|3pbOb;9NI;O z*q^wYK5Xof!OOSoO*i;`c@;_CE8PxEwVvwM^_1R1Ni@h&CGv>Lf$aKf-EJN@y5KxF z8kErol!O!$8h9N!036gpDvv&n!k5@K;fDRdp;g1F>{T$;kIf~fyh(g}2jKXcP=uKf zZXV-8Ut2RO*o@{YTe0TZuJ*D9E53uzrLyQ_r=4pQT+^t`2xIWPxM9TUMg41@>B!v2 z?R!)VQR~!Cio;T34tfVw^f7uq9@p-`$adz2#g|=5+|vk}1%wz8UvLVDQbQ^v6KSyg zXTQbviiY(*uPDGfVv%f#X+nWC`PMjbZ~!7Dxjcd&du6)uw&8l$rTPK$>M%4h+1uXp za8A=@EcUh5V>V>^W9V~YN2!?$Y^?Dgy~h3o>miNU(vkN0?FNGQ?{E}t&>yIe>YNaX zpqxE#Pnoq-r1f|0C}S%NB@6i%{_CfEvfCfX@ASsBcy1%& zRlDxtrhlFqcWPlDtd|eV?%p^yc=51Cj8Uu|G^t@cXwc*M&%%KjNq=*|$uIy)R29NV z)C=4h9(2LG_S01vVghh(xbdAim35Vl6#$Jaz zAgWy-O+O)6`4_UqsYS|__j5o~m?{C;l9+Y#a61FqKbP(cM=TZ~!zp4E-%7gKIwo<4 zddnj%YoAGaJHHD@MF^kWerAWN$&sJ09(H-l@6Ie8RFxgiz7TRi`w%GvXxkaWB!3c6 z0)_l##p^~$tCCLI#4v(vMLDrgHZQ6YR1ZfQ0ywHF)fQONR_@OuP87Cd&`IbRXb{Kx zJR)s<8P5FNF6V0lbWTgZ z7;$a+EGR*XpkUE)yX#U$ffyM3I4V=T!wceqUsa_?asme#Nz!kH`mu!3M}bJJ-y^N; zvVC%#Dd9jQz4gHvUEj{@v^zVFoBLebXFGl8l$bM%2LRvatMBDrIro&2QeF=m*!Cp& zRtIDd)UM!zJU`)ctpxc53PUpi;;Bd25rdtA*i^~Ds`NTlffunAjvGJWm(kB2KsfPQ zopqxhVw4s$!P$=t^ft@-nRA70yd6ERu8aY}KeEKwlp}6XZiA?}{m*CU%epk1(!G9u;U>JgOmZL>~nuE$YV&$Erzj zhw^WGD>t=rLfEpQ_Ujg0Nx)g+v(-Z{{o$dHmfrLFvnluuR*Y!M_cyRl1Rt6>8)@@t z<-7jbe|+(ZM^Q@*jR(Iyw3LOG5;z0!KE_34&5M?w_8@qOTmU z*Piy=w&tn!-@N$i?>9a4arr$pIY}>$Y{rNRd>s6-&f+sufM^j!@_Y}LQ+U@^eDR7+ zI291rJW3kEj#y*^5COXg5Ea9MU{bZjfv+g@pL{tGEsqUTp>|qOB4vW3EM$lWN-$b6 z)Lbl&@&1(db!q7Op$^0i$zE0WuS-HU;1AQw_Rj)YM9id&8x2!7O$Tx$Hfe augmentationMap = {37, "crop"}, {38, "crop_mirror_normalize"}, {39, "resize_crop_mirror"}, + {45, "color_temperature"}, {49, "box_filter"}, {54, "gaussian_filter"}, {63, "phase"},