Half-inl.h 9.1 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182183184185186187188189190191192193194195196197198199200201202203204205206207208209210211212213214215216217218219220221222223224225226227228229230231232233234235236237238239240241242243244245246247248249250251252253254255256257258259260261262263264265266267268269270271272273274275276277278279280281282283284285286287288289290291292293294295296297298299300301302303304305306307308309310311312313314315316317318319320321322323324325326
  1. #pragma once
  2. #include <c10/macros/Macros.h>
  3. #include <cstring>
  4. #include <limits>
  5. #ifdef __CUDACC__
  6. #include <cuda_fp16.h>
  7. #endif
  8. #ifdef __HIPCC__
  9. #include <hip/hip_fp16.h>
  10. #endif
  11. #if defined(CL_SYCL_LANGUAGE_VERSION)
  12. #include <CL/sycl.hpp> // for SYCL 1.2.1
  13. #elif defined(SYCL_LANGUAGE_VERSION)
  14. #include <sycl/sycl.hpp> // for SYCL 2020
  15. #endif
  16. C10_CLANG_DIAGNOSTIC_PUSH()
  17. #if C10_CLANG_HAS_WARNING("-Wimplicit-int-float-conversion")
  18. C10_CLANG_DIAGNOSTIC_IGNORE("-Wimplicit-int-float-conversion")
  19. #endif
  20. namespace c10 {
  21. /// Constructors
  22. inline C10_HOST_DEVICE Half::Half(float value)
  23. :
  24. #if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__)
  25. x(__half_as_short(__float2half(value)))
  26. #elif defined(__SYCL_DEVICE_ONLY__)
  27. x(sycl::bit_cast<uint16_t>(sycl::half(value)))
  28. #else
  29. x(detail::fp16_ieee_from_fp32_value(value))
  30. #endif
  31. {
  32. }
  33. /// Implicit conversions
  34. inline C10_HOST_DEVICE Half::operator float() const {
  35. #if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__)
  36. return __half2float(*reinterpret_cast<const __half*>(&x));
  37. #elif defined(__SYCL_DEVICE_ONLY__)
  38. return float(sycl::bit_cast<sycl::half>(x));
  39. #else
  40. return detail::fp16_ieee_to_fp32_value(x);
  41. #endif
  42. }
  43. #if defined(__CUDACC__) || defined(__HIPCC__)
  44. inline C10_HOST_DEVICE Half::Half(const __half& value) {
  45. x = *reinterpret_cast<const unsigned short*>(&value);
  46. }
  47. inline C10_HOST_DEVICE Half::operator __half() const {
  48. return *reinterpret_cast<const __half*>(&x);
  49. }
  50. #endif
  51. #ifdef SYCL_LANGUAGE_VERSION
  52. inline C10_HOST_DEVICE Half::Half(const sycl::half& value) {
  53. x = *reinterpret_cast<const unsigned short*>(&value);
  54. }
  55. inline C10_HOST_DEVICE Half::operator sycl::half() const {
  56. return *reinterpret_cast<const sycl::half*>(&x);
  57. }
  58. #endif
  59. // CUDA intrinsics
  60. #if (defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 350)) || \
  61. (defined(__clang__) && defined(__CUDA__))
  62. inline __device__ Half __ldg(const Half* ptr) {
  63. return __ldg(reinterpret_cast<const __half*>(ptr));
  64. }
  65. #endif
  66. /// Arithmetic
  67. inline C10_HOST_DEVICE Half operator+(const Half& a, const Half& b) {
  68. return static_cast<float>(a) + static_cast<float>(b);
  69. }
  70. inline C10_HOST_DEVICE Half operator-(const Half& a, const Half& b) {
  71. return static_cast<float>(a) - static_cast<float>(b);
  72. }
  73. inline C10_HOST_DEVICE Half operator*(const Half& a, const Half& b) {
  74. return static_cast<float>(a) * static_cast<float>(b);
  75. }
  76. inline C10_HOST_DEVICE Half operator/(const Half& a, const Half& b)
  77. __ubsan_ignore_float_divide_by_zero__ {
  78. return static_cast<float>(a) / static_cast<float>(b);
  79. }
  80. inline C10_HOST_DEVICE Half operator-(const Half& a) {
  81. #if (defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 530) || \
  82. defined(__HIP_DEVICE_COMPILE__)
  83. return __hneg(a);
  84. #elif defined(__SYCL_DEVICE_ONLY__)
  85. return -sycl::bit_cast<sycl::half>(a);
  86. #else
  87. return -static_cast<float>(a);
  88. #endif
  89. }
  90. inline C10_HOST_DEVICE Half& operator+=(Half& a, const Half& b) {
  91. a = a + b;
  92. return a;
  93. }
  94. inline C10_HOST_DEVICE Half& operator-=(Half& a, const Half& b) {
  95. a = a - b;
  96. return a;
  97. }
  98. inline C10_HOST_DEVICE Half& operator*=(Half& a, const Half& b) {
  99. a = a * b;
  100. return a;
  101. }
  102. inline C10_HOST_DEVICE Half& operator/=(Half& a, const Half& b) {
  103. a = a / b;
  104. return a;
  105. }
  106. /// Arithmetic with floats
  107. inline C10_HOST_DEVICE float operator+(Half a, float b) {
  108. return static_cast<float>(a) + b;
  109. }
  110. inline C10_HOST_DEVICE float operator-(Half a, float b) {
  111. return static_cast<float>(a) - b;
  112. }
  113. inline C10_HOST_DEVICE float operator*(Half a, float b) {
  114. return static_cast<float>(a) * b;
  115. }
  116. inline C10_HOST_DEVICE float operator/(Half a, float b)
  117. __ubsan_ignore_float_divide_by_zero__ {
  118. return static_cast<float>(a) / b;
  119. }
  120. inline C10_HOST_DEVICE float operator+(float a, Half b) {
  121. return a + static_cast<float>(b);
  122. }
  123. inline C10_HOST_DEVICE float operator-(float a, Half b) {
  124. return a - static_cast<float>(b);
  125. }
  126. inline C10_HOST_DEVICE float operator*(float a, Half b) {
  127. return a * static_cast<float>(b);
  128. }
  129. inline C10_HOST_DEVICE float operator/(float a, Half b)
  130. __ubsan_ignore_float_divide_by_zero__ {
  131. return a / static_cast<float>(b);
  132. }
  133. inline C10_HOST_DEVICE float& operator+=(float& a, const Half& b) {
  134. return a += static_cast<float>(b);
  135. }
  136. inline C10_HOST_DEVICE float& operator-=(float& a, const Half& b) {
  137. return a -= static_cast<float>(b);
  138. }
  139. inline C10_HOST_DEVICE float& operator*=(float& a, const Half& b) {
  140. return a *= static_cast<float>(b);
  141. }
  142. inline C10_HOST_DEVICE float& operator/=(float& a, const Half& b) {
  143. return a /= static_cast<float>(b);
  144. }
  145. /// Arithmetic with doubles
  146. inline C10_HOST_DEVICE double operator+(Half a, double b) {
  147. return static_cast<double>(a) + b;
  148. }
  149. inline C10_HOST_DEVICE double operator-(Half a, double b) {
  150. return static_cast<double>(a) - b;
  151. }
  152. inline C10_HOST_DEVICE double operator*(Half a, double b) {
  153. return static_cast<double>(a) * b;
  154. }
  155. inline C10_HOST_DEVICE double operator/(Half a, double b)
  156. __ubsan_ignore_float_divide_by_zero__ {
  157. return static_cast<double>(a) / b;
  158. }
  159. inline C10_HOST_DEVICE double operator+(double a, Half b) {
  160. return a + static_cast<double>(b);
  161. }
  162. inline C10_HOST_DEVICE double operator-(double a, Half b) {
  163. return a - static_cast<double>(b);
  164. }
  165. inline C10_HOST_DEVICE double operator*(double a, Half b) {
  166. return a * static_cast<double>(b);
  167. }
  168. inline C10_HOST_DEVICE double operator/(double a, Half b)
  169. __ubsan_ignore_float_divide_by_zero__ {
  170. return a / static_cast<double>(b);
  171. }
  172. /// Arithmetic with ints
  173. inline C10_HOST_DEVICE Half operator+(Half a, int b) {
  174. return a + static_cast<Half>(b);
  175. }
  176. inline C10_HOST_DEVICE Half operator-(Half a, int b) {
  177. return a - static_cast<Half>(b);
  178. }
  179. inline C10_HOST_DEVICE Half operator*(Half a, int b) {
  180. return a * static_cast<Half>(b);
  181. }
  182. inline C10_HOST_DEVICE Half operator/(Half a, int b) {
  183. return a / static_cast<Half>(b);
  184. }
  185. inline C10_HOST_DEVICE Half operator+(int a, Half b) {
  186. return static_cast<Half>(a) + b;
  187. }
  188. inline C10_HOST_DEVICE Half operator-(int a, Half b) {
  189. return static_cast<Half>(a) - b;
  190. }
  191. inline C10_HOST_DEVICE Half operator*(int a, Half b) {
  192. return static_cast<Half>(a) * b;
  193. }
  194. inline C10_HOST_DEVICE Half operator/(int a, Half b) {
  195. return static_cast<Half>(a) / b;
  196. }
  197. //// Arithmetic with int64_t
  198. inline C10_HOST_DEVICE Half operator+(Half a, int64_t b) {
  199. return a + static_cast<Half>(b);
  200. }
  201. inline C10_HOST_DEVICE Half operator-(Half a, int64_t b) {
  202. return a - static_cast<Half>(b);
  203. }
  204. inline C10_HOST_DEVICE Half operator*(Half a, int64_t b) {
  205. return a * static_cast<Half>(b);
  206. }
  207. inline C10_HOST_DEVICE Half operator/(Half a, int64_t b) {
  208. return a / static_cast<Half>(b);
  209. }
  210. inline C10_HOST_DEVICE Half operator+(int64_t a, Half b) {
  211. return static_cast<Half>(a) + b;
  212. }
  213. inline C10_HOST_DEVICE Half operator-(int64_t a, Half b) {
  214. return static_cast<Half>(a) - b;
  215. }
  216. inline C10_HOST_DEVICE Half operator*(int64_t a, Half b) {
  217. return static_cast<Half>(a) * b;
  218. }
  219. inline C10_HOST_DEVICE Half operator/(int64_t a, Half b) {
  220. return static_cast<Half>(a) / b;
  221. }
  222. /// NOTE: we do not define comparisons directly and instead rely on the implicit
  223. /// conversion from c10::Half to float.
  224. } // namespace c10
  225. namespace std {
  226. template <>
  227. class numeric_limits<c10::Half> {
  228. public:
  229. static constexpr bool is_specialized = true;
  230. static constexpr bool is_signed = true;
  231. static constexpr bool is_integer = false;
  232. static constexpr bool is_exact = false;
  233. static constexpr bool has_infinity = true;
  234. static constexpr bool has_quiet_NaN = true;
  235. static constexpr bool has_signaling_NaN = true;
  236. static constexpr auto has_denorm = numeric_limits<float>::has_denorm;
  237. static constexpr auto has_denorm_loss =
  238. numeric_limits<float>::has_denorm_loss;
  239. static constexpr auto round_style = numeric_limits<float>::round_style;
  240. static constexpr bool is_iec559 = true;
  241. static constexpr bool is_bounded = true;
  242. static constexpr bool is_modulo = false;
  243. static constexpr int digits = 11;
  244. static constexpr int digits10 = 3;
  245. static constexpr int max_digits10 = 5;
  246. static constexpr int radix = 2;
  247. static constexpr int min_exponent = -13;
  248. static constexpr int min_exponent10 = -4;
  249. static constexpr int max_exponent = 16;
  250. static constexpr int max_exponent10 = 4;
  251. static constexpr auto traps = numeric_limits<float>::traps;
  252. static constexpr auto tinyness_before =
  253. numeric_limits<float>::tinyness_before;
  254. static constexpr c10::Half min() {
  255. return c10::Half(0x0400, c10::Half::from_bits());
  256. }
  257. static constexpr c10::Half lowest() {
  258. return c10::Half(0xFBFF, c10::Half::from_bits());
  259. }
  260. static constexpr c10::Half max() {
  261. return c10::Half(0x7BFF, c10::Half::from_bits());
  262. }
  263. static constexpr c10::Half epsilon() {
  264. return c10::Half(0x1400, c10::Half::from_bits());
  265. }
  266. static constexpr c10::Half round_error() {
  267. return c10::Half(0x3800, c10::Half::from_bits());
  268. }
  269. static constexpr c10::Half infinity() {
  270. return c10::Half(0x7C00, c10::Half::from_bits());
  271. }
  272. static constexpr c10::Half quiet_NaN() {
  273. return c10::Half(0x7E00, c10::Half::from_bits());
  274. }
  275. static constexpr c10::Half signaling_NaN() {
  276. return c10::Half(0x7D00, c10::Half::from_bits());
  277. }
  278. static constexpr c10::Half denorm_min() {
  279. return c10::Half(0x0001, c10::Half::from_bits());
  280. }
  281. };
  282. } // namespace std
  283. C10_CLANG_DIAGNOSTIC_POP()