/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-roccv/checkouts/latest/include/kernels/device/gamma_contrast_device.hpp Source File

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-roccv/checkouts/latest/include/kernels/device/gamma_contrast_device.hpp Source File#

2 min read time

Applies to Linux

rocCV: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-roccv/checkouts/latest/include/kernels/device/gamma_contrast_device.hpp Source File

Go to the documentation of this file.

  
 #pragma once
  
 #include <hip/hip_runtime.h>
 #include "core/detail/math/vectorized_type_math.hpp"
 #include "core/detail/casting.hpp"
  
 #include "operator_types.h"
  
 namespace Kernels {
 namespace Device {
 template <typename SrcWrapper, typename DstWrapper>
 __global__ void gamma_contrast(SrcWrapper input, DstWrapper output, float gamma) {
     using namespace roccv::detail;  // For RangeCast, NumElements, etc.
     using src_type = typename SrcWrapper::ValueType;
     using dst_type = typename DstWrapper::ValueType;
     using work_type = MakeType<float, NumElements<src_type>>;
     
     const int x = threadIdx.x + blockIdx.x * blockDim.x;
     const int y = threadIdx.y + blockIdx.y * blockDim.y;
     const int batch = blockIdx.z;
  
     if (x >= output.width() || y >= output.height() || batch >= output.batches()) return;
     
     auto inVal = (RangeCast<work_type>(input.at(batch, y, x, 0)));
     work_type result = math::vpowf(inVal, gamma);
     if constexpr (NumElements<dst_type> == 4) {
         output.at(batch, y, x, 0) = RangeCast<dst_type>((MakeType<float, 4>){result.x, result.y, result.z, inVal.w});
     } else {
         output.at(batch, y, x, 0) = RangeCast<dst_type>(result);
     }
     
 }
 }  // namespace Device
 }  // namespace Kernels