/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-roccv/checkouts/latest/include/core/detail/casting.hpp Source File

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-roccv/checkouts/latest/include/core/detail/casting.hpp Source File#

7 min read time

Applies to Linux

rocCV: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-roccv/checkouts/latest/include/core/detail/casting.hpp Source File
casting.hpp
Go to the documentation of this file.
1 /*
2  * Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved.
3  * Permission is hereby granted, free of charge, to any person obtaining a copy
4  * of this software and associated documentation files (the "Software"), to deal
5  * in the Software without restriction, including without limitation the rights
6  * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
7  * copies of the Software, and to permit persons to whom the Software is
8  * furnished to do so, subject to the following conditions:
9  *
10  * The above copyright notice and this permission notice shall be included in
11  * all copies or substantial portions of the Software.
12  *
13  * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
14  * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
15  * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
16  * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
17  * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
18  * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
19  * THE SOFTWARE.
20  */
21 
22 #pragma once
23 
24 #include <algorithm>
25 
26 #include "core/detail/type_traits.hpp"
27 
28 namespace roccv::detail {
29 
33 template <typename T, typename U, class = std::enable_if_t<!IsCompound<T> && !IsCompound<U>>>
34 __device__ __host__ T ScalarSaturateCast(U v) {
35  constexpr bool smallToBig = sizeof(U) <= sizeof(T);
36  constexpr bool bigToSmall = !smallToBig;
37 
38  if constexpr (std::is_integral_v<T> && std::is_floating_point_v<U>) {
39  // Any float -> any integral
40  return static_cast<T>(std::clamp<U>(std::round(v), static_cast<U>(std::numeric_limits<T>::min()),
41  static_cast<U>(std::numeric_limits<T>::max())));
42  } else if constexpr (std::is_integral_v<T> && std::is_integral_v<U> && std::is_signed_v<U> && std::is_signed_v<T> &&
43  smallToBig) {
44  // Any integral signed -> Any integral unsigned, small -> big or equal
45  return v <= 0 ? 0 : static_cast<T>(v);
46  } else if constexpr (std::is_integral_v<U> && std::is_integral_v<T> &&
47  ((std::is_signed_v<U> && std::is_signed_v<T>) ||
48  (std::is_unsigned_v<U> && std::is_unsigned_v<T>)) &&
49  bigToSmall) {
50  // Any integral signed -> Any integral signed, big -> small
51  // Any integral unsigned -> Any integral unsigned, big -> small
52  return v <= std::numeric_limits<T>::min()
53  ? std::numeric_limits<T>::min()
54  : (v >= std::numeric_limits<T>::max() ? std::numeric_limits<T>::max() : static_cast<T>(v));
55  } else if constexpr (std::is_integral_v<U> && std::is_unsigned_v<U> && std::is_integral_v<T> &&
56  std::is_signed_v<T>) {
57  // Any integral unsigned -> Any integral signed, small -> big or equal
58  return v >= std::numeric_limits<T>::max() ? std::numeric_limits<T>::max() : static_cast<T>(v);
59  } else if constexpr (std::is_integral_v<U> && std::is_signed_v<U> && std::is_integral_v<T> &&
60  std::is_unsigned_v<T> && bigToSmall) {
61  // Any integral signed -> Any integral unsigned, big -> small
62  return v <= static_cast<U>(std::numeric_limits<T>::min())
63  ? std::numeric_limits<T>::min()
64  : (v >= static_cast<U>(std::numeric_limits<T>::max()) ? std::numeric_limits<T>::max
65  : static_cast<T>(v));
66  } else {
67  // All other cases fall into this
68  return v;
69  }
70 }
71 
83 template <typename T, typename U,
84  class = std::enable_if_t<(HasTypeTraits<T> && HasTypeTraits<U>) && (NumElements<T> <= NumElements<U>)>>
85 __device__ __host__ T SaturateCast(U v) {
86  if constexpr (std::is_same_v<T, U>) {
87  return v;
88  }
89 
90  T ret{};
91 
92  GetElement(ret, 0) = ScalarSaturateCast<BaseType<T>>(GetElement(v, 0));
93  if constexpr (NumElements<T> >= 2) GetElement(ret, 1) = ScalarSaturateCast<BaseType<T>>(GetElement(v, 1));
94  if constexpr (NumElements<T> >= 3) GetElement(ret, 2) = ScalarSaturateCast<BaseType<T>>(GetElement(v, 2));
95  if constexpr (NumElements<T> >= 4) GetElement(ret, 3) = ScalarSaturateCast<BaseType<T>>(GetElement(v, 3));
96 
97  return ret;
98 }
99 
103 template <typename T, typename U,
104  class = std::enable_if_t<(HasTypeTraits<T> && HasTypeTraits<U>) && (!IsCompound<T> && !IsCompound<U>)>>
105 __device__ __host__ T ScalarRangeCast(U v) {
106  if constexpr (std::is_same_v<T, U>) {
107  // Types are the same, no work needed
108  return v;
109  }
110 
111  else if constexpr (std::is_integral_v<T> && std::is_floating_point_v<U> && std::is_signed_v<T>) {
112  // Float to signed integers
113  return v >= T{1} ? std::numeric_limits<T>::max()
114  : v <= T{-1} ? std::numeric_limits<T>::min()
115  : static_cast<T>(std::round(static_cast<U>(std::numeric_limits<T>::max()) * v));
116  }
117 
118  else if constexpr (std::is_integral_v<T> && std::is_floating_point_v<U> && std::is_unsigned_v<T>) {
119  // float to unsigned integers
120  return v >= T{1} ? std::numeric_limits<T>::max()
121  : v <= T{0} ? 0
122  : static_cast<T>(lrintf(static_cast<U>(std::numeric_limits<T>::max()) * v));
123  }
124 
125  else if constexpr (std::is_floating_point_v<T> && std::is_integral_v<U> && std::is_signed_v<U>) {
126  // Signed integer to float
127  constexpr T invmax = T{1} / static_cast<T>(std::numeric_limits<U>::max());
128  T out = static_cast<T>(v) * invmax;
129  return out < T{-1} ? T{-1} : out;
130  }
131 
132  else if constexpr (std::is_floating_point_v<T> && std::is_integral_v<U> && std::is_unsigned_v<U>) {
133  // Unsigned integer to float
134  constexpr T invmax = T{1} / static_cast<T>(std::numeric_limits<U>::max());
135  return static_cast<T>(v) * invmax;
136  }
137 
138  else {
139  // All other cases reduce to a saturate cast
140  return ScalarSaturateCast<T>(v);
141  }
142 }
143 
162 template <typename T, typename U,
163  class = std::enable_if_t<(HasTypeTraits<T> && HasTypeTraits<U>) && NumElements<T> <= NumElements<U>>>
164 __device__ __host__ T RangeCast(U v) {
165  if constexpr (std::is_same_v<T, U>) {
166  return v;
167  }
168 
169  T ret{};
170 
171  GetElement(ret, 0) = ScalarRangeCast<BaseType<T>>(GetElement(v, 0));
172  if constexpr (NumElements<T> >= 2) GetElement(ret, 1) = ScalarRangeCast<BaseType<T>>(GetElement(v, 1));
173  if constexpr (NumElements<T> >= 3) GetElement(ret, 2) = ScalarRangeCast<BaseType<T>>(GetElement(v, 2));
174  if constexpr (NumElements<T> >= 4) GetElement(ret, 3) = ScalarRangeCast<BaseType<T>>(GetElement(v, 3));
175 
176  return ret;
177 }
178 
187 template <typename T, typename U,
188  class = std::enable_if_t<(HasTypeTraits<T> && HasTypeTraits<U>) && NumElements<T> <= NumElements<U>>>
189 __device__ __host__ T StaticCast(U v) {
190  if constexpr (std::is_same_v<T, U>) {
191  // Both same type, just return the value.
192  return v;
193  } else if constexpr (!IsCompound<T> && !IsCompound<U>) {
194  // Both scalar values. Reduces to a standard static cast.
195  return static_cast<T>(v);
196  } else {
197  // Vector types. Perform casting on each element.
198  T ret{};
199  GetElement(ret, 0) = StaticCast<BaseType<T>>(GetElement(v, 0));
200  if constexpr (NumElements<T> >= 2) GetElement(ret, 1) = StaticCast<BaseType<T>>(GetElement(v, 1));
201  if constexpr (NumElements<T> >= 3) GetElement(ret, 2) = StaticCast<BaseType<T>>(GetElement(v, 2));
202  if constexpr (NumElements<T> >= 4) GetElement(ret, 3) = StaticCast<BaseType<T>>(GetElement(v, 3));
203 
204  return ret;
205  }
206 }
207 } // namespace roccv::detail
roccv::detail
Definition: strided_data_wrap.hpp:33
roccv::detail::SaturateCast
__device__ __host__ T SaturateCast(U v)
Performs a saturation cast from one type to another. Each type must have type traits supported....
Definition: casting.hpp:85
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::ScalarSaturateCast
__device__ __host__ T ScalarSaturateCast(U v)
ScalarSaturateCast is for implementation purposes only. Use SaturateCast directly.
Definition: casting.hpp:34
roccv::detail::GetElement
__host__ __device__ RT & GetElement(T &v, int idx)
Definition: type_traits.hpp:128
roccv::detail::ScalarRangeCast
__device__ __host__ T ScalarRangeCast(U v)
ScalarRangeCast is for implementation purposes only. Use RangeCast directly instead.
Definition: casting.hpp:105
roccv::detail::NumElements
constexpr int NumElements
Returns the number of elements in a HIP vectorized type. For example: uchar3 will return 3,...
Definition: type_traits.hpp:94
roccv::detail::HasTypeTraits
constexpr bool HasTypeTraits
Returns whether the datatype passed in has type traits associated with it.
Definition: type_traits.hpp:125
type_traits.hpp