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

5 min read time

Applies to Linux

• include
• kernels
• host

thresholding_host.hpp

Go to the documentation of this file.

1 

23 #pragma once

24 

25 #include <hip/hip_runtime.h>

26 #include <core/wrappers/generic_tensor_wrapper.hpp>

27 #include "core/detail/type_traits.hpp"

28 #include "core/detail/casting.hpp"

29 #include "operator_types.h"

30 

31 namespace Kernels {

32 namespace Host {

33 template <typename SrcWrapper, typename DstWrapper>

34 void binary_generic(SrcWrapper input, DstWrapper output, roccv::GenericTensorWrapper<double> thresh,

35  roccv::GenericTensorWrapper<double> maxVal, const int32_t maxBatchSize) {

36  using namespace roccv::detail;

37  using src_type = typename SrcWrapper::ValueType;

38  using dst_type = typename DstWrapper::ValueType;

39  using base_type = BaseType<dst_type>;

40 

41  for (int z_idx = 0; z_idx < maxBatchSize; z_idx++) {

42  double th = thresh.at(z_idx);

43  double mv = maxVal.at(z_idx);

44  for (int y_idx = 0; y_idx < output.height(); y_idx++) {

45  for (int x_idx = 0; x_idx < output.width(); x_idx++) {

46  src_type inputVal = input.at(z_idx, y_idx, x_idx, 0);

47  dst_type outputVal;

48  for (int i = 0; i < output.channels(); i++) {

49  double ip = StaticCast<double>(GetElement(inputVal, i));

50  double outVal = ip > th ? mv : 0;

51  GetElement(outputVal, i) = StaticCast<base_type>(outVal);

52  }

53  output.at(z_idx, y_idx, x_idx, 0) = outputVal;

54  }

55  }

56  }

57 }

58 

59 template <typename SrcWrapper, typename DstWrapper>

60 void binary_inv_generic(SrcWrapper input, DstWrapper output, roccv::GenericTensorWrapper<double> thresh,

61  roccv::GenericTensorWrapper<double> maxVal, const int32_t maxBatchSize) {

62  using namespace roccv::detail;

63  using src_type = typename SrcWrapper::ValueType;

64  using dst_type = typename DstWrapper::ValueType;

65  using base_type = BaseType<dst_type>;

66 

67  for (int z_idx = 0; z_idx < maxBatchSize; z_idx++) {

68  double th = thresh.at(z_idx);

69  double mv = maxVal.at(z_idx);

70  for (int y_idx = 0; y_idx < output.height(); y_idx++) {

71  for (int x_idx = 0; x_idx < output.width(); x_idx++) {

72  src_type inputVal = input.at(z_idx, y_idx, x_idx, 0);

73  dst_type outputVal;

74  for (int i = 0; i < output.channels(); i++) {

75  double ip = StaticCast<double>(GetElement(inputVal, i));

76  double outVal = ip > th ? 0 : mv;

77  GetElement(outputVal, i) = StaticCast<base_type>(outVal);

78  }

79  output.at(z_idx, y_idx, x_idx, 0) = outputVal;

80  }

81  }

82  }

83 }

84 

85 template <typename SrcWrapper, typename DstWrapper>

86 void trunc_generic(SrcWrapper input, DstWrapper output,

87  roccv::GenericTensorWrapper<double> thresh, const int32_t maxBatchSize) {

88  using namespace roccv::detail;

89  using src_type = typename SrcWrapper::ValueType;

90  using dst_type = typename DstWrapper::ValueType;

91  using base_type = BaseType<dst_type>;

92 

93  for (int z_idx = 0; z_idx < maxBatchSize; z_idx++) {

94  double th = thresh.at(z_idx);

95  for (int y_idx = 0; y_idx < output.height(); y_idx++) {

96  for (int x_idx = 0; x_idx < output.width(); x_idx++) {

97  src_type inputVal = input.at(z_idx, y_idx, x_idx, 0);

98  dst_type outputVal;

99  for (int i = 0; i < output.channels(); i++) {

100  double ip = StaticCast<double>(GetElement(inputVal, i));

101  double outVal = ip > th ? th : ip;

102  GetElement(outputVal, i) = StaticCast<base_type>(outVal);

103  }

104  output.at(z_idx, y_idx, x_idx, 0) = outputVal;

105  }

106  }

107  }

108 }

109 

110 template <typename SrcWrapper, typename DstWrapper>

111 void tozero_generic(SrcWrapper input, DstWrapper output,

112  roccv::GenericTensorWrapper<double> thresh, const int32_t maxBatchSize) {

113  using namespace roccv::detail;

114  using src_type = typename SrcWrapper::ValueType;

115  using dst_type = typename DstWrapper::ValueType;

116  using base_type = BaseType<dst_type>;

117 

118  for (int z_idx = 0; z_idx < maxBatchSize; z_idx++) {

119  double th = thresh.at(z_idx);

120  for (int y_idx = 0; y_idx < output.height(); y_idx++) {

121  for (int x_idx = 0; x_idx < output.width(); x_idx++) {

122  src_type inputVal = input.at(z_idx, y_idx, x_idx, 0);

123  dst_type outputVal;

124  for (int i = 0; i < output.channels(); i++) {

125  double ip = StaticCast<double>(GetElement(inputVal, i));

126  double outVal = ip > th ? ip : 0;

127  GetElement(outputVal, i) = StaticCast<base_type>(outVal);

128  }

129  output.at(z_idx, y_idx, x_idx, 0) = outputVal;

130  }

131  }

132  }

133 }

134 

135 template <typename SrcWrapper, typename DstWrapper>

136 void tozeroinv_generic(SrcWrapper input, DstWrapper output,

137  roccv::GenericTensorWrapper<double> thresh, const int32_t maxBatchSize) {

138  using namespace roccv::detail;

139  using src_type = typename SrcWrapper::ValueType;

140  using dst_type = typename DstWrapper::ValueType;

141  using base_type = BaseType<dst_type>;

142 

143  for (int z_idx = 0; z_idx < maxBatchSize; z_idx++) {

144  double th = thresh.at(z_idx);

145  for (int y_idx = 0; y_idx < output.height(); y_idx++) {

146  for (int x_idx = 0; x_idx < output.width(); x_idx++) {

147  src_type inputVal = input.at(z_idx, y_idx, x_idx, 0);

148  dst_type outputVal;

149  for (int i = 0; i < output.channels(); i++) {

150  double ip = StaticCast<double>(GetElement(inputVal, i));

151  double outVal = ip > th ? 0 : ip;

152  GetElement(outputVal, i) = StaticCast<base_type>(outVal);

153  }

154  output.at(z_idx, y_idx, x_idx, 0) = outputVal;

155  }

156  }

157  }

158 }

159 } // namespace Host

160 } // namespace Kernels

casting.hpp

roccv::GenericTensorWrapper

Definition: generic_tensor_wrapper.hpp:28

roccv::GenericTensorWrapper::at

__device__ __host__ T & at(ARGS... idx)

Definition: generic_tensor_wrapper.hpp:48

generic_tensor_wrapper.hpp

Kernels::Host::tozero_generic

void tozero_generic(SrcWrapper input, DstWrapper output, roccv::GenericTensorWrapper< double > thresh, const int32_t maxBatchSize)

Definition: thresholding_host.hpp:111

Kernels::Host::binary_generic

void binary_generic(SrcWrapper input, DstWrapper output, roccv::GenericTensorWrapper< double > thresh, roccv::GenericTensorWrapper< double > maxVal, const int32_t maxBatchSize)

Definition: thresholding_host.hpp:34

Kernels::Host::trunc_generic

void trunc_generic(SrcWrapper input, DstWrapper output, roccv::GenericTensorWrapper< double > thresh, const int32_t maxBatchSize)

Definition: thresholding_host.hpp:86

Kernels::Host::binary_inv_generic

void binary_inv_generic(SrcWrapper input, DstWrapper output, roccv::GenericTensorWrapper< double > thresh, roccv::GenericTensorWrapper< double > maxVal, const int32_t maxBatchSize)

Definition: thresholding_host.hpp:60

Kernels::Host::tozeroinv_generic

void tozeroinv_generic(SrcWrapper input, DstWrapper output, roccv::GenericTensorWrapper< double > thresh, const int32_t maxBatchSize)

Definition: thresholding_host.hpp:136

Kernels

Definition: non_max_suppression_helpers.hpp:26

roccv::detail

Definition: strided_data_wrap.hpp:33

roccv::detail::BaseType

typename TypeTraits< T >::base_type BaseType

Returns the base type of a given HIP vectorized type.

Definition: type_traits.hpp:117

roccv::detail::GetElement

__host__ __device__ RT & GetElement(T &v, int idx)

Definition: type_traits.hpp:128

operator_types.h

type_traits.hpp