GPU/PacketMath.h
Go to the documentation of this file.
1 // This file is part of Eigen, a lightweight C++ template library
2 // for linear algebra.
3 //
4 // Copyright (C) 2014 Benoit Steiner <benoit.steiner.goog@gmail.com>
5 //
6 // This Source Code Form is subject to the terms of the Mozilla
7 // Public License v. 2.0. If a copy of the MPL was not distributed
8 // with this file, You can obtain one at http://mozilla.org/MPL/2.0/.
9 
10 #ifndef EIGEN_PACKET_MATH_GPU_H
11 #define EIGEN_PACKET_MATH_GPU_H
12 
13 namespace Eigen {
14 
15 namespace internal {
16 
17 // Read-only data cached load available.
18 #if defined(EIGEN_HIP_DEVICE_COMPILE) || (defined(EIGEN_CUDA_ARCH) && EIGEN_CUDA_ARCH >= 350)
19 #define EIGEN_GPU_HAS_LDG 1
20 #endif
21 
22 // FP16 math available.
23 #if (defined(EIGEN_CUDA_ARCH) && EIGEN_CUDA_ARCH >= 530)
24 #define EIGEN_CUDA_HAS_FP16_ARITHMETIC 1
25 #endif
26 
27 #if defined(EIGEN_HIP_DEVICE_COMPILE) || defined(EIGEN_CUDA_HAS_FP16_ARITHMETIC)
28 #define EIGEN_GPU_HAS_FP16_ARITHMETIC 1
29 #endif
30 
31 // Make sure this is only available when targeting a GPU: we don't want to
32 // introduce conflicts between these packet_traits definitions and the ones
33 // we'll use on the host side (SSE, AVX, ...)
34 #if defined(EIGEN_GPUCC) && defined(EIGEN_USE_GPU)
35 
36 template<> struct is_arithmetic<float4> { enum { value = true }; };
37 template<> struct is_arithmetic<double2> { enum { value = true }; };
38 
39 template<> struct packet_traits<float> : default_packet_traits
40 {
41  typedef float4 type;
42  typedef float4 half;
43  enum {
44  Vectorizable = 1,
45  AlignedOnScalar = 1,
46  size=4,
47  HasHalfPacket = 0,
48 
49  HasDiv = 1,
50  HasSin = 0,
51  HasCos = 0,
52  HasLog = 1,
53  HasExp = 1,
54  HasSqrt = 1,
55  HasRsqrt = 1,
56  HasLGamma = 1,
57  HasDiGamma = 1,
58  HasZeta = 1,
59  HasPolygamma = 1,
60  HasErf = 1,
61  HasErfc = 1,
62  HasNdtri = 1,
63  HasBessel = 1,
64  HasIGamma = 1,
65  HasIGammaDerA = 1,
67  HasIGammac = 1,
68  HasBetaInc = 1,
69 
70  HasBlend = 0,
71  HasFloor = 1,
72  };
73 };
74 
75 template<> struct packet_traits<double> : default_packet_traits
76 {
77  typedef double2 type;
78  typedef double2 half;
79  enum {
80  Vectorizable = 1,
81  AlignedOnScalar = 1,
82  size=2,
83  HasHalfPacket = 0,
84 
85  HasDiv = 1,
86  HasLog = 1,
87  HasExp = 1,
88  HasSqrt = 1,
89  HasRsqrt = 1,
90  HasLGamma = 1,
91  HasDiGamma = 1,
92  HasZeta = 1,
93  HasPolygamma = 1,
94  HasErf = 1,
95  HasErfc = 1,
96  HasNdtri = 1,
97  HasBessel = 1,
98  HasIGamma = 1,
99  HasIGammaDerA = 1,
101  HasIGammac = 1,
102  HasBetaInc = 1,
103 
104  HasBlend = 0,
105  HasFloor = 1,
106  };
107 };
108 
109 
110 template<> struct unpacket_traits<float4> { typedef float type; enum {size=4, alignment=Aligned16, vectorizable=true, masked_load_available=false, masked_store_available=false}; typedef float4 half; };
111 template<> struct unpacket_traits<double2> { typedef double type; enum {size=2, alignment=Aligned16, vectorizable=true, masked_load_available=false, masked_store_available=false}; typedef double2 half; };
112 
113 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pset1<float4>(const float& from) {
114  return make_float4(from, from, from, from);
115 }
116 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 pset1<double2>(const double& from) {
117  return make_double2(from, from);
118 }
119 
120 // We need to distinguish ‘clang as the CUDA compiler’ from ‘clang as the host compiler,
121 // invoked by NVCC’ (e.g. on MacOS). The former needs to see both host and device implementation
122 // of the functions, while the latter can only deal with one of them.
123 #if defined(EIGEN_CUDA_ARCH) || defined(EIGEN_HIPCC) || (defined(EIGEN_CUDACC) && EIGEN_COMP_CLANG && !EIGEN_COMP_NVCC)
124 namespace {
125 
126 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float bitwise_and(const float& a,
127  const float& b) {
128  return __int_as_float(__float_as_int(a) & __float_as_int(b));
129 }
130 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double bitwise_and(const double& a,
131  const double& b) {
132  return __longlong_as_double(__double_as_longlong(a) &
133  __double_as_longlong(b));
134 }
135 
136 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float bitwise_or(const float& a,
137  const float& b) {
138  return __int_as_float(__float_as_int(a) | __float_as_int(b));
139 }
140 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double bitwise_or(const double& a,
141  const double& b) {
142  return __longlong_as_double(__double_as_longlong(a) |
143  __double_as_longlong(b));
144 }
145 
146 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float bitwise_xor(const float& a,
147  const float& b) {
148  return __int_as_float(__float_as_int(a) ^ __float_as_int(b));
149 }
150 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double bitwise_xor(const double& a,
151  const double& b) {
152  return __longlong_as_double(__double_as_longlong(a) ^
153  __double_as_longlong(b));
154 }
155 
156 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float bitwise_andnot(const float& a,
157  const float& b) {
158  return __int_as_float(__float_as_int(a) & ~__float_as_int(b));
159 }
160 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double bitwise_andnot(const double& a,
161  const double& b) {
162  return __longlong_as_double(__double_as_longlong(a) &
163  ~__double_as_longlong(b));
164 }
165 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float eq_mask(const float& a,
166  const float& b) {
167  return __int_as_float(a == b ? 0xffffffffu : 0u);
168 }
169 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double eq_mask(const double& a,
170  const double& b) {
171  return __longlong_as_double(a == b ? 0xffffffffffffffffull : 0ull);
172 }
173 
174 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float lt_mask(const float& a,
175  const float& b) {
176  return __int_as_float(a < b ? 0xffffffffu : 0u);
177 }
178 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double lt_mask(const double& a,
179  const double& b) {
180  return __longlong_as_double(a < b ? 0xffffffffffffffffull : 0ull);
181 }
182 
183 } // namespace
184 
185 template <>
186 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pand<float4>(const float4& a,
187  const float4& b) {
188  return make_float4(bitwise_and(a.x, b.x), bitwise_and(a.y, b.y),
189  bitwise_and(a.z, b.z), bitwise_and(a.w, b.w));
190 }
191 template <>
193  const double2& b) {
194  return make_double2(bitwise_and(a.x, b.x), bitwise_and(a.y, b.y));
195 }
196 
197 template <>
198 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 por<float4>(const float4& a,
199  const float4& b) {
200  return make_float4(bitwise_or(a.x, b.x), bitwise_or(a.y, b.y),
201  bitwise_or(a.z, b.z), bitwise_or(a.w, b.w));
202 }
203 template <>
205  const double2& b) {
206  return make_double2(bitwise_or(a.x, b.x), bitwise_or(a.y, b.y));
207 }
208 
209 template <>
210 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pxor<float4>(const float4& a,
211  const float4& b) {
212  return make_float4(bitwise_xor(a.x, b.x), bitwise_xor(a.y, b.y),
213  bitwise_xor(a.z, b.z), bitwise_xor(a.w, b.w));
214 }
215 template <>
217  const double2& b) {
218  return make_double2(bitwise_xor(a.x, b.x), bitwise_xor(a.y, b.y));
219 }
220 
221 template <>
222 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pandnot<float4>(const float4& a,
223  const float4& b) {
224  return make_float4(bitwise_andnot(a.x, b.x), bitwise_andnot(a.y, b.y),
225  bitwise_andnot(a.z, b.z), bitwise_andnot(a.w, b.w));
226 }
227 template <>
229 pandnot<double2>(const double2& a, const double2& b) {
230  return make_double2(bitwise_andnot(a.x, b.x), bitwise_andnot(a.y, b.y));
231 }
232 
233 template <>
234 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pcmp_eq<float4>(const float4& a,
235  const float4& b) {
236  return make_float4(eq_mask(a.x, b.x), eq_mask(a.y, b.y), eq_mask(a.z, b.z),
237  eq_mask(a.w, b.w));
238 }
239 template <>
240 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pcmp_lt<float4>(const float4& a,
241  const float4& b) {
242  return make_float4(lt_mask(a.x, b.x), lt_mask(a.y, b.y), lt_mask(a.z, b.z),
243  lt_mask(a.w, b.w));
244 }
245 template <>
247 pcmp_eq<double2>(const double2& a, const double2& b) {
248  return make_double2(eq_mask(a.x, b.x), eq_mask(a.y, b.y));
249 }
250 template <>
252 pcmp_lt<double2>(const double2& a, const double2& b) {
253  return make_double2(lt_mask(a.x, b.x), lt_mask(a.y, b.y));
254 }
255 #endif // defined(EIGEN_CUDA_ARCH) || defined(EIGEN_HIPCC) || (defined(EIGEN_CUDACC) && EIGEN_COMP_CLANG && !EIGEN_COMP_NVCC)
256 
257 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 plset<float4>(const float& a) {
258  return make_float4(a, a+1, a+2, a+3);
259 }
260 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 plset<double2>(const double& a) {
261  return make_double2(a, a+1);
262 }
263 
264 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 padd<float4>(const float4& a, const float4& b) {
265  return make_float4(a.x+b.x, a.y+b.y, a.z+b.z, a.w+b.w);
266 }
267 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 padd<double2>(const double2& a, const double2& b) {
268  return make_double2(a.x+b.x, a.y+b.y);
269 }
270 
271 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 psub<float4>(const float4& a, const float4& b) {
272  return make_float4(a.x-b.x, a.y-b.y, a.z-b.z, a.w-b.w);
273 }
274 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 psub<double2>(const double2& a, const double2& b) {
275  return make_double2(a.x-b.x, a.y-b.y);
276 }
277 
278 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pnegate(const float4& a) {
279  return make_float4(-a.x, -a.y, -a.z, -a.w);
280 }
282  return make_double2(-a.x, -a.y);
283 }
284 
285 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pconj(const float4& a) { return a; }
286 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 pconj(const double2& a) { return a; }
287 
288 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pmul<float4>(const float4& a, const float4& b) {
289  return make_float4(a.x*b.x, a.y*b.y, a.z*b.z, a.w*b.w);
290 }
291 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 pmul<double2>(const double2& a, const double2& b) {
292  return make_double2(a.x*b.x, a.y*b.y);
293 }
294 
295 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pdiv<float4>(const float4& a, const float4& b) {
296  return make_float4(a.x/b.x, a.y/b.y, a.z/b.z, a.w/b.w);
297 }
298 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 pdiv<double2>(const double2& a, const double2& b) {
299  return make_double2(a.x/b.x, a.y/b.y);
300 }
301 
302 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pmin<float4>(const float4& a, const float4& b) {
303  return make_float4(fminf(a.x, b.x), fminf(a.y, b.y), fminf(a.z, b.z), fminf(a.w, b.w));
304 }
305 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 pmin<double2>(const double2& a, const double2& b) {
306  return make_double2(fmin(a.x, b.x), fmin(a.y, b.y));
307 }
308 
309 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pmax<float4>(const float4& a, const float4& b) {
310  return make_float4(fmaxf(a.x, b.x), fmaxf(a.y, b.y), fmaxf(a.z, b.z), fmaxf(a.w, b.w));
311 }
312 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 pmax<double2>(const double2& a, const double2& b) {
313  return make_double2(fmax(a.x, b.x), fmax(a.y, b.y));
314 }
315 
316 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pload<float4>(const float* from) {
317  return *reinterpret_cast<const float4*>(from);
318 }
319 
320 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 pload<double2>(const double* from) {
321  return *reinterpret_cast<const double2*>(from);
322 }
323 
324 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 ploadu<float4>(const float* from) {
325  return make_float4(from[0], from[1], from[2], from[3]);
326 }
327 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 ploadu<double2>(const double* from) {
328  return make_double2(from[0], from[1]);
329 }
330 
331 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 ploaddup<float4>(const float* from) {
332  return make_float4(from[0], from[0], from[1], from[1]);
333 }
334 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 ploaddup<double2>(const double* from) {
335  return make_double2(from[0], from[0]);
336 }
337 
338 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void pstore<float>(float* to, const float4& from) {
339  *reinterpret_cast<float4*>(to) = from;
340 }
341 
342 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void pstore<double>(double* to, const double2& from) {
343  *reinterpret_cast<double2*>(to) = from;
344 }
345 
346 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void pstoreu<float>(float* to, const float4& from) {
347  to[0] = from.x;
348  to[1] = from.y;
349  to[2] = from.z;
350  to[3] = from.w;
351 }
352 
353 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void pstoreu<double>(double* to, const double2& from) {
354  to[0] = from.x;
355  to[1] = from.y;
356 }
357 
358 template<>
359 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float4 ploadt_ro<float4, Aligned>(const float* from) {
360 #if defined(EIGEN_GPU_HAS_LDG)
361  return __ldg((const float4*)from);
362 #else
363  return make_float4(from[0], from[1], from[2], from[3]);
364 #endif
365 }
366 template<>
367 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double2 ploadt_ro<double2, Aligned>(const double* from) {
368 #if defined(EIGEN_GPU_HAS_LDG)
369  return __ldg((const double2*)from);
370 #else
371  return make_double2(from[0], from[1]);
372 #endif
373 }
374 
375 template<>
376 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float4 ploadt_ro<float4, Unaligned>(const float* from) {
377 #if defined(EIGEN_GPU_HAS_LDG)
378  return make_float4(__ldg(from+0), __ldg(from+1), __ldg(from+2), __ldg(from+3));
379 #else
380  return make_float4(from[0], from[1], from[2], from[3]);
381 #endif
382 }
383 template<>
384 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double2 ploadt_ro<double2, Unaligned>(const double* from) {
385 #if defined(EIGEN_GPU_HAS_LDG)
386  return make_double2(__ldg(from+0), __ldg(from+1));
387 #else
388  return make_double2(from[0], from[1]);
389 #endif
390 }
391 
392 template<> EIGEN_DEVICE_FUNC inline float4 pgather<float, float4>(const float* from, Index stride) {
393  return make_float4(from[0*stride], from[1*stride], from[2*stride], from[3*stride]);
394 }
395 
396 template<> EIGEN_DEVICE_FUNC inline double2 pgather<double, double2>(const double* from, Index stride) {
397  return make_double2(from[0*stride], from[1*stride]);
398 }
399 
400 template<> EIGEN_DEVICE_FUNC inline void pscatter<float, float4>(float* to, const float4& from, Index stride) {
401  to[stride*0] = from.x;
402  to[stride*1] = from.y;
403  to[stride*2] = from.z;
404  to[stride*3] = from.w;
405 }
406 template<> EIGEN_DEVICE_FUNC inline void pscatter<double, double2>(double* to, const double2& from, Index stride) {
407  to[stride*0] = from.x;
408  to[stride*1] = from.y;
409 }
410 
411 template<> EIGEN_DEVICE_FUNC inline float pfirst<float4>(const float4& a) {
412  return a.x;
413 }
414 template<> EIGEN_DEVICE_FUNC inline double pfirst<double2>(const double2& a) {
415  return a.x;
416 }
417 
418 template<> EIGEN_DEVICE_FUNC inline float predux<float4>(const float4& a) {
419  return a.x + a.y + a.z + a.w;
420 }
421 template<> EIGEN_DEVICE_FUNC inline double predux<double2>(const double2& a) {
422  return a.x + a.y;
423 }
424 
425 template<> EIGEN_DEVICE_FUNC inline float predux_max<float4>(const float4& a) {
426  return fmaxf(fmaxf(a.x, a.y), fmaxf(a.z, a.w));
427 }
428 template<> EIGEN_DEVICE_FUNC inline double predux_max<double2>(const double2& a) {
429  return fmax(a.x, a.y);
430 }
431 
432 template<> EIGEN_DEVICE_FUNC inline float predux_min<float4>(const float4& a) {
433  return fminf(fminf(a.x, a.y), fminf(a.z, a.w));
434 }
435 template<> EIGEN_DEVICE_FUNC inline double predux_min<double2>(const double2& a) {
436  return fmin(a.x, a.y);
437 }
438 
439 template<> EIGEN_DEVICE_FUNC inline float predux_mul<float4>(const float4& a) {
440  return a.x * a.y * a.z * a.w;
441 }
442 template<> EIGEN_DEVICE_FUNC inline double predux_mul<double2>(const double2& a) {
443  return a.x * a.y;
444 }
445 
446 template<> EIGEN_DEVICE_FUNC inline float4 pabs<float4>(const float4& a) {
447  return make_float4(fabsf(a.x), fabsf(a.y), fabsf(a.z), fabsf(a.w));
448 }
449 template<> EIGEN_DEVICE_FUNC inline double2 pabs<double2>(const double2& a) {
450  return make_double2(fabs(a.x), fabs(a.y));
451 }
452 
453 template<> EIGEN_DEVICE_FUNC inline float4 pfloor<float4>(const float4& a) {
454  return make_float4(floorf(a.x), floorf(a.y), floorf(a.z), floorf(a.w));
455 }
456 template<> EIGEN_DEVICE_FUNC inline double2 pfloor<double2>(const double2& a) {
457  return make_double2(floor(a.x), floor(a.y));
458 }
459 
460 EIGEN_DEVICE_FUNC inline void
461 ptranspose(PacketBlock<float4,4>& kernel) {
462  float tmp = kernel.packet[0].y;
463  kernel.packet[0].y = kernel.packet[1].x;
464  kernel.packet[1].x = tmp;
465 
466  tmp = kernel.packet[0].z;
467  kernel.packet[0].z = kernel.packet[2].x;
468  kernel.packet[2].x = tmp;
469 
470  tmp = kernel.packet[0].w;
471  kernel.packet[0].w = kernel.packet[3].x;
472  kernel.packet[3].x = tmp;
473 
474  tmp = kernel.packet[1].z;
475  kernel.packet[1].z = kernel.packet[2].y;
476  kernel.packet[2].y = tmp;
477 
478  tmp = kernel.packet[1].w;
479  kernel.packet[1].w = kernel.packet[3].y;
480  kernel.packet[3].y = tmp;
481 
482  tmp = kernel.packet[2].w;
483  kernel.packet[2].w = kernel.packet[3].z;
484  kernel.packet[3].z = tmp;
485 }
486 
487 EIGEN_DEVICE_FUNC inline void
488 ptranspose(PacketBlock<double2,2>& kernel) {
489  double tmp = kernel.packet[0].y;
490  kernel.packet[0].y = kernel.packet[1].x;
491  kernel.packet[1].x = tmp;
492 }
493 
494 #endif // defined(EIGEN_GPUCC) && defined(EIGEN_USE_GPU)
495 
496 // Packet4h2 must be defined in the macro without EIGEN_CUDA_ARCH, meaning
497 // its corresponding packet_traits<Eigen::half> must be visible on host.
498 #if defined(EIGEN_HAS_CUDA_FP16) || defined(EIGEN_HAS_HIP_FP16)
499 
500 typedef ulonglong2 Packet4h2;
501 template<> struct unpacket_traits<Packet4h2> { typedef Eigen::half type; enum {size=8, alignment=Aligned16, vectorizable=true, masked_load_available=false, masked_store_available=false}; typedef Packet4h2 half; };
502 template<> struct is_arithmetic<Packet4h2> { enum { value = true }; };
503 
504 template<> struct unpacket_traits<half2> { typedef Eigen::half type; enum {size=2, alignment=Aligned16, vectorizable=true, masked_load_available=false, masked_store_available=false}; typedef half2 half; };
505 template<> struct is_arithmetic<half2> { enum { value = true }; };
506 
507 template<> struct packet_traits<Eigen::half> : default_packet_traits
508 {
509  typedef Packet4h2 type;
510  typedef Packet4h2 half;
511  enum {
512  Vectorizable = 1,
513  AlignedOnScalar = 1,
514  size=8,
515  HasHalfPacket = 0,
516  HasAdd = 1,
517  HasSub = 1,
518  HasMul = 1,
519  HasDiv = 1,
520  HasSqrt = 1,
521  HasRsqrt = 1,
522  HasExp = 1,
523  HasExpm1 = 1,
524  HasLog = 1,
525  HasLog1p = 1
526  };
527 };
528 
529 namespace {
530 // This is equivalent to make_half2, which is undocumented and doesn't seem to always exist.
531 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 combine_half(const __half& a, const __half& b) {
532 #if defined(EIGEN_GPU_COMPILE_PHASE)
533  return __halves2half2(a, b);
534 #else
535  // Round-about way since __halves2half2 is a __device__ function.
536  return __floats2half2_rn(__half2float(a), __half2float(b));
537 #endif
538 }
539 
540 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE __half get_half2_low(const half2& a) {
541 #if defined(EIGEN_GPU_COMPILE_PHASE)
542  return __low2half(a);
543 #else
544  return __float2half(__low2float(a));
545 #endif
546 }
547 
548 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE __half get_half2_high(const half2& a) {
549 #if defined(EIGEN_GPU_COMPILE_PHASE)
550  return __high2half(a);
551 #else
552  return __float2half(__high2float(a));
553 #endif
554 }
555 } // namespace
556 
557 template<>
558 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pset1<half2>(const Eigen::half& from) {
559 #if defined(EIGEN_GPU_COMPILE_PHASE)
560  return __half2half2(from);
561 #else
562  const float f = __half2float(from);
563  return __floats2half2_rn(f, f);
564 #endif
565 }
566 
567 template <>
569 pset1<Packet4h2>(const Eigen::half& from) {
570  Packet4h2 r;
571  half2* p_alias = reinterpret_cast<half2*>(&r);
572  p_alias[0] = pset1<half2>(from);
573  p_alias[1] = pset1<half2>(from);
574  p_alias[2] = pset1<half2>(from);
575  p_alias[3] = pset1<half2>(from);
576  return r;
577 }
578 
579 // We now need this visible on both host and device.
580 // #if defined(EIGEN_CUDA_ARCH) || defined(EIGEN_HIPCC) || (defined(EIGEN_CUDACC) && EIGEN_COMP_CLANG && !EIGEN_COMP_NVCC)
581 namespace {
582 
584  return *reinterpret_cast<const half2*>(from);
585 }
586 
588  return combine_half(from[0], from[1]);
589 }
590 
592  return combine_half(from[0], from[0]);
593 }
594 
596  const half2& from) {
597  *reinterpret_cast<half2*>(to) = from;
598 }
599 
601  const half2& from) {
602  to[0] = get_half2_low(from);
603  to[1] = get_half2_high(from);
604 }
605 
606 
607 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE half2 ploadt_ro_aligned(
608  const Eigen::half* from) {
609 #if defined(EIGEN_GPU_HAS_LDG)
610  // Input is guaranteed to be properly aligned.
611  return __ldg(reinterpret_cast<const half2*>(from));
612 #else
613  return combine_half(*(from+0), *(from+1));
614 #endif
615 }
616 
617 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE half2 ploadt_ro_unaligned(
618  const Eigen::half* from) {
619 #if defined(EIGEN_GPU_HAS_LDG)
620  return __halves2half2(__ldg(from+0), __ldg(from+1));
621 #else
622  return combine_half(*(from+0), *(from+1));
623 #endif
624 }
625 
627  Index stride) {
628  return combine_half(from[0*stride], from[1*stride]);
629 }
630 
632  Eigen::half* to, const half2& from, Index stride) {
633  to[stride*0] = get_half2_low(from);
634  to[stride*1] = get_half2_high(from);
635 }
636 
638  return get_half2_low(a);
639 }
640 
641 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pabs(const half2& a) {
642  half a1 = get_half2_low(a);
643  half a2 = get_half2_high(a);
644  half result1 = half_impl::raw_uint16_to_half(a1.x & 0x7FFF);
645  half result2 = half_impl::raw_uint16_to_half(a2.x & 0x7FFF);
646  return combine_half(result1, result2);
647 }
648 
649 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 ptrue(const half2& /*a*/) {
650  half true_half = half_impl::raw_uint16_to_half(0xffffu);
651  return pset1<half2>(true_half);
652 }
653 
654 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pzero(const half2& /*a*/) {
655  half false_half = half_impl::raw_uint16_to_half(0x0000u);
656  return pset1<half2>(false_half);
657 }
658 
660 ptranspose(PacketBlock<half2,2>& kernel) {
661  __half a1 = get_half2_low(kernel.packet[0]);
662  __half a2 = get_half2_high(kernel.packet[0]);
663  __half b1 = get_half2_low(kernel.packet[1]);
664  __half b2 = get_half2_high(kernel.packet[1]);
665  kernel.packet[0] = combine_half(a1, b1);
666  kernel.packet[1] = combine_half(a2, b2);
667 }
668 
670 #if defined(EIGEN_GPU_HAS_FP16_ARITHMETIC)
671  return __halves2half2(a, __hadd(a, __float2half(1.0f)));
672 #else
673  float f = __half2float(a) + 1.0f;
674  return combine_half(a, __float2half(f));
675 #endif
676 }
677 
678 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pselect(const half2& mask,
679  const half2& a,
680  const half2& b) {
681  half mask_low = get_half2_low(mask);
682  half mask_high = get_half2_high(mask);
683  half result_low = mask_low == half(0) ? get_half2_low(b) : get_half2_low(a);
684  half result_high = mask_high == half(0) ? get_half2_high(b) : get_half2_high(a);
685  return combine_half(result_low, result_high);
686 }
687 
689  const half2& b) {
690  half true_half = half_impl::raw_uint16_to_half(0xffffu);
691  half false_half = half_impl::raw_uint16_to_half(0x0000u);
692  half a1 = get_half2_low(a);
693  half a2 = get_half2_high(a);
694  half b1 = get_half2_low(b);
695  half b2 = get_half2_high(b);
696  half eq1 = __half2float(a1) == __half2float(b1) ? true_half : false_half;
697  half eq2 = __half2float(a2) == __half2float(b2) ? true_half : false_half;
698  return combine_half(eq1, eq2);
699 }
700 
702  const half2& b) {
703  half true_half = half_impl::raw_uint16_to_half(0xffffu);
704  half false_half = half_impl::raw_uint16_to_half(0x0000u);
705  half a1 = get_half2_low(a);
706  half a2 = get_half2_high(a);
707  half b1 = get_half2_low(b);
708  half b2 = get_half2_high(b);
709  half eq1 = __half2float(a1) < __half2float(b1) ? true_half : false_half;
710  half eq2 = __half2float(a2) < __half2float(b2) ? true_half : false_half;
711  return combine_half(eq1, eq2);
712 }
713 
714 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pand(const half2& a,
715  const half2& b) {
716  half a1 = get_half2_low(a);
717  half a2 = get_half2_high(a);
718  half b1 = get_half2_low(b);
719  half b2 = get_half2_high(b);
720  half result1 = half_impl::raw_uint16_to_half(a1.x & b1.x);
721  half result2 = half_impl::raw_uint16_to_half(a2.x & b2.x);
722  return combine_half(result1, result2);
723 }
724 
725 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 por(const half2& a,
726  const half2& b) {
727  half a1 = get_half2_low(a);
728  half a2 = get_half2_high(a);
729  half b1 = get_half2_low(b);
730  half b2 = get_half2_high(b);
731  half result1 = half_impl::raw_uint16_to_half(a1.x | b1.x);
732  half result2 = half_impl::raw_uint16_to_half(a2.x | b2.x);
733  return combine_half(result1, result2);
734 }
735 
736 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pxor(const half2& a,
737  const half2& b) {
738  half a1 = get_half2_low(a);
739  half a2 = get_half2_high(a);
740  half b1 = get_half2_low(b);
741  half b2 = get_half2_high(b);
742  half result1 = half_impl::raw_uint16_to_half(a1.x ^ b1.x);
743  half result2 = half_impl::raw_uint16_to_half(a2.x ^ b2.x);
744  return combine_half(result1, result2);
745 }
746 
748  const half2& b) {
749  half a1 = get_half2_low(a);
750  half a2 = get_half2_high(a);
751  half b1 = get_half2_low(b);
752  half b2 = get_half2_high(b);
753  half result1 = half_impl::raw_uint16_to_half(a1.x & ~b1.x);
754  half result2 = half_impl::raw_uint16_to_half(a2.x & ~b2.x);
755  return combine_half(result1, result2);
756 }
757 
758 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 padd(const half2& a,
759  const half2& b) {
760 #if defined(EIGEN_GPU_HAS_FP16_ARITHMETIC)
761  return __hadd2(a, b);
762 #else
763  float a1 = __low2float(a);
764  float a2 = __high2float(a);
765  float b1 = __low2float(b);
766  float b2 = __high2float(b);
767  float r1 = a1 + b1;
768  float r2 = a2 + b2;
769  return __floats2half2_rn(r1, r2);
770 #endif
771 }
772 
773 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 psub(const half2& a,
774  const half2& b) {
775 #if defined(EIGEN_GPU_HAS_FP16_ARITHMETIC)
776  return __hsub2(a, b);
777 #else
778  float a1 = __low2float(a);
779  float a2 = __high2float(a);
780  float b1 = __low2float(b);
781  float b2 = __high2float(b);
782  float r1 = a1 - b1;
783  float r2 = a2 - b2;
784  return __floats2half2_rn(r1, r2);
785 #endif
786 }
787 
788 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pnegate(const half2& a) {
789 #if defined(EIGEN_GPU_HAS_FP16_ARITHMETIC)
790  return __hneg2(a);
791 #else
792  float a1 = __low2float(a);
793  float a2 = __high2float(a);
794  return __floats2half2_rn(-a1, -a2);
795 #endif
796 }
797 
798 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pconj(const half2& a) { return a; }
799 
800 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pmul(const half2& a,
801  const half2& b) {
802 #if defined(EIGEN_GPU_HAS_FP16_ARITHMETIC)
803  return __hmul2(a, b);
804 #else
805  float a1 = __low2float(a);
806  float a2 = __high2float(a);
807  float b1 = __low2float(b);
808  float b2 = __high2float(b);
809  float r1 = a1 * b1;
810  float r2 = a2 * b2;
811  return __floats2half2_rn(r1, r2);
812 #endif
813 }
814 
815 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pmadd(const half2& a,
816  const half2& b,
817  const half2& c) {
818 #if defined(EIGEN_GPU_HAS_FP16_ARITHMETIC)
819  return __hfma2(a, b, c);
820 #else
821  float a1 = __low2float(a);
822  float a2 = __high2float(a);
823  float b1 = __low2float(b);
824  float b2 = __high2float(b);
825  float c1 = __low2float(c);
826  float c2 = __high2float(c);
827  float r1 = a1 * b1 + c1;
828  float r2 = a2 * b2 + c2;
829  return __floats2half2_rn(r1, r2);
830 #endif
831 }
832 
833 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pdiv(const half2& a,
834  const half2& b) {
835 #if defined(EIGEN_GPU_HAS_FP16_ARITHMETIC)
836  return __h2div(a, b);
837 #else
838  float a1 = __low2float(a);
839  float a2 = __high2float(a);
840  float b1 = __low2float(b);
841  float b2 = __high2float(b);
842  float r1 = a1 / b1;
843  float r2 = a2 / b2;
844  return __floats2half2_rn(r1, r2);
845 #endif
846 }
847 
848 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pmin(const half2& a,
849  const half2& b) {
850  float a1 = __low2float(a);
851  float a2 = __high2float(a);
852  float b1 = __low2float(b);
853  float b2 = __high2float(b);
854  __half r1 = a1 < b1 ? get_half2_low(a) : get_half2_low(b);
855  __half r2 = a2 < b2 ? get_half2_high(a) : get_half2_high(b);
856  return combine_half(r1, r2);
857 }
858 
859 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pmax(const half2& a,
860  const half2& b) {
861  float a1 = __low2float(a);
862  float a2 = __high2float(a);
863  float b1 = __low2float(b);
864  float b2 = __high2float(b);
865  __half r1 = a1 > b1 ? get_half2_low(a) : get_half2_low(b);
866  __half r2 = a2 > b2 ? get_half2_high(a) : get_half2_high(b);
867  return combine_half(r1, r2);
868 }
869 
871 #if defined(EIGEN_GPU_HAS_FP16_ARITHMETIC)
872  return __hadd(__low2half(a), __high2half(a));
873 #else
874  float a1 = __low2float(a);
875  float a2 = __high2float(a);
876  return Eigen::half(__float2half(a1 + a2));
877 #endif
878 }
879 
881 #if defined(EIGEN_GPU_HAS_FP16_ARITHMETIC)
882  __half first = __low2half(a);
883  __half second = __high2half(a);
884  return __hgt(first, second) ? first : second;
885 #else
886  float a1 = __low2float(a);
887  float a2 = __high2float(a);
888  return a1 > a2 ? get_half2_low(a) : get_half2_high(a);
889 #endif
890 }
891 
893 #if defined(EIGEN_GPU_HAS_FP16_ARITHMETIC)
894  __half first = __low2half(a);
895  __half second = __high2half(a);
896  return __hlt(first, second) ? first : second;
897 #else
898  float a1 = __low2float(a);
899  float a2 = __high2float(a);
900  return a1 < a2 ? get_half2_low(a) : get_half2_high(a);
901 #endif
902 }
903 
905 #if defined(EIGEN_GPU_HAS_FP16_ARITHMETIC)
906  return __hmul(__low2half(a), __high2half(a));
907 #else
908  float a1 = __low2float(a);
909  float a2 = __high2float(a);
910  return Eigen::half(__float2half(a1 * a2));
911 #endif
912 }
913 
914 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 plog1p(const half2& a) {
915  float a1 = __low2float(a);
916  float a2 = __high2float(a);
917  float r1 = log1pf(a1);
918  float r2 = log1pf(a2);
919  return __floats2half2_rn(r1, r2);
920 }
921 
922 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pexpm1(const half2& a) {
923  float a1 = __low2float(a);
924  float a2 = __high2float(a);
925  float r1 = expm1f(a1);
926  float r2 = expm1f(a2);
927  return __floats2half2_rn(r1, r2);
928 }
929 
930 #if (EIGEN_CUDA_SDK_VER >= 80000 && defined(EIGEN_CUDA_HAS_FP16_ARITHMETIC)) || \
931  defined(EIGEN_HIP_DEVICE_COMPILE)
932 
934 half2 plog(const half2& a) {
935  return h2log(a);
936 }
937 
939 half2 pexp(const half2& a) {
940  return h2exp(a);
941 }
942 
944 half2 psqrt(const half2& a) {
945  return h2sqrt(a);
946 }
947 
949 half2 prsqrt(const half2& a) {
950  return h2rsqrt(a);
951 }
952 
953 #else
954 
955 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 plog(const half2& a) {
956  float a1 = __low2float(a);
957  float a2 = __high2float(a);
958  float r1 = logf(a1);
959  float r2 = logf(a2);
960  return __floats2half2_rn(r1, r2);
961 }
962 
963 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pexp(const half2& a) {
964  float a1 = __low2float(a);
965  float a2 = __high2float(a);
966  float r1 = expf(a1);
967  float r2 = expf(a2);
968  return __floats2half2_rn(r1, r2);
969 }
970 
971 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 psqrt(const half2& a) {
972  float a1 = __low2float(a);
973  float a2 = __high2float(a);
974  float r1 = sqrtf(a1);
975  float r2 = sqrtf(a2);
976  return __floats2half2_rn(r1, r2);
977 }
978 
979 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 prsqrt(const half2& a) {
980  float a1 = __low2float(a);
981  float a2 = __high2float(a);
982  float r1 = rsqrtf(a1);
983  float r2 = rsqrtf(a2);
984  return __floats2half2_rn(r1, r2);
985 }
986 #endif
987 } // namespace
988 
989 template <>
991 pload<Packet4h2>(const Eigen::half* from) {
992  return *reinterpret_cast<const Packet4h2*>(from);
993 }
994 
995 // unaligned load;
996 template <>
998 ploadu<Packet4h2>(const Eigen::half* from) {
999  Packet4h2 r;
1000  half2* p_alias = reinterpret_cast<half2*>(&r);
1001  p_alias[0] = ploadu(from + 0);
1002  p_alias[1] = ploadu(from + 2);
1003  p_alias[2] = ploadu(from + 4);
1004  p_alias[3] = ploadu(from + 6);
1005  return r;
1006 }
1007 
1008 template <>
1010 ploaddup<Packet4h2>(const Eigen::half* from) {
1011  Packet4h2 r;
1012  half2* p_alias = reinterpret_cast<half2*>(&r);
1013  p_alias[0] = ploaddup(from + 0);
1014  p_alias[1] = ploaddup(from + 1);
1015  p_alias[2] = ploaddup(from + 2);
1016  p_alias[3] = ploaddup(from + 3);
1017  return r;
1018 }
1019 
1020 template <>
1021 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void pstore<Eigen::half>(
1022  Eigen::half* to, const Packet4h2& from) {
1023  *reinterpret_cast<Packet4h2*>(to) = from;
1024 }
1025 
1026 template <>
1027 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void pstoreu<Eigen::half>(
1028  Eigen::half* to, const Packet4h2& from) {
1029  const half2* from_alias = reinterpret_cast<const half2*>(&from);
1030  pstoreu(to + 0,from_alias[0]);
1031  pstoreu(to + 2,from_alias[1]);
1032  pstoreu(to + 4,from_alias[2]);
1033  pstoreu(to + 6,from_alias[3]);
1034 }
1035 
1036 template <>
1038 ploadt_ro<Packet4h2, Aligned>(const Eigen::half* from) {
1039 #if defined(EIGEN_GPU_HAS_LDG)
1040  Packet4h2 r;
1041  r = __ldg(reinterpret_cast<const Packet4h2*>(from));
1042  return r;
1043 #else
1044  Packet4h2 r;
1045  half2* r_alias = reinterpret_cast<half2*>(&r);
1046  r_alias[0] = ploadt_ro_aligned(from + 0);
1047  r_alias[1] = ploadt_ro_aligned(from + 2);
1048  r_alias[2] = ploadt_ro_aligned(from + 4);
1049  r_alias[3] = ploadt_ro_aligned(from + 6);
1050  return r;
1051 #endif
1052 }
1053 
1054 template <>
1056 ploadt_ro<Packet4h2, Unaligned>(const Eigen::half* from) {
1057  Packet4h2 r;
1058  half2* r_alias = reinterpret_cast<half2*>(&r);
1059  r_alias[0] = ploadt_ro_unaligned(from + 0);
1060  r_alias[1] = ploadt_ro_unaligned(from + 2);
1061  r_alias[2] = ploadt_ro_unaligned(from + 4);
1062  r_alias[3] = ploadt_ro_unaligned(from + 6);
1063  return r;
1064 }
1065 
1066 template <>
1068 pgather<Eigen::half, Packet4h2>(const Eigen::half* from, Index stride) {
1069  Packet4h2 r;
1070  half2* p_alias = reinterpret_cast<half2*>(&r);
1071  p_alias[0] = combine_half(from[0 * stride], from[1 * stride]);
1072  p_alias[1] = combine_half(from[2 * stride], from[3 * stride]);
1073  p_alias[2] = combine_half(from[4 * stride], from[5 * stride]);
1074  p_alias[3] = combine_half(from[6 * stride], from[7 * stride]);
1075  return r;
1076 }
1077 
1078 template <>
1079 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void pscatter<Eigen::half, Packet4h2>(
1080  Eigen::half* to, const Packet4h2& from, Index stride) {
1081  const half2* from_alias = reinterpret_cast<const half2*>(&from);
1082  pscatter(to + stride * 0, from_alias[0], stride);
1083  pscatter(to + stride * 2, from_alias[1], stride);
1084  pscatter(to + stride * 4, from_alias[2], stride);
1085  pscatter(to + stride * 6, from_alias[3], stride);
1086 }
1087 
1088 template <>
1090  const Packet4h2& a) {
1091  return pfirst(*(reinterpret_cast<const half2*>(&a)));
1092 }
1093 
1094 template <>
1095 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 pabs<Packet4h2>(
1096  const Packet4h2& a) {
1097  Packet4h2 r;
1098  half2* p_alias = reinterpret_cast<half2*>(&r);
1099  const half2* a_alias = reinterpret_cast<const half2*>(&a);
1100  p_alias[0] = pabs(a_alias[0]);
1101  p_alias[1] = pabs(a_alias[1]);
1102  p_alias[2] = pabs(a_alias[2]);
1103  p_alias[3] = pabs(a_alias[3]);
1104  return r;
1105 }
1106 
1107 template <>
1108 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 ptrue<Packet4h2>(
1109  const Packet4h2& /*a*/) {
1110  half true_half = half_impl::raw_uint16_to_half(0xffffu);
1111  return pset1<Packet4h2>(true_half);
1112 }
1113 
1114 template <>
1115 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 pzero<Packet4h2>(const Packet4h2& /*a*/) {
1116  half false_half = half_impl::raw_uint16_to_half(0x0000u);
1117  return pset1<Packet4h2>(false_half);
1118 }
1119 
1120 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void ptranspose_double(
1121  double* d_row0, double* d_row1, double* d_row2, double* d_row3,
1122  double* d_row4, double* d_row5, double* d_row6, double* d_row7) {
1123  double d_tmp;
1124  d_tmp = d_row0[1];
1125  d_row0[1] = d_row4[0];
1126  d_row4[0] = d_tmp;
1127 
1128  d_tmp = d_row1[1];
1129  d_row1[1] = d_row5[0];
1130  d_row5[0] = d_tmp;
1131 
1132  d_tmp = d_row2[1];
1133  d_row2[1] = d_row6[0];
1134  d_row6[0] = d_tmp;
1135 
1136  d_tmp = d_row3[1];
1137  d_row3[1] = d_row7[0];
1138  d_row7[0] = d_tmp;
1139 }
1140 
1141 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void ptranspose_half2(
1142  half2* f_row0, half2* f_row1, half2* f_row2, half2* f_row3) {
1143  half2 f_tmp;
1144  f_tmp = f_row0[1];
1145  f_row0[1] = f_row2[0];
1146  f_row2[0] = f_tmp;
1147 
1148  f_tmp = f_row1[1];
1149  f_row1[1] = f_row3[0];
1150  f_row3[0] = f_tmp;
1151 }
1152 
1154 ptranspose_half(half2& f0, half2& f1) {
1155  __half a1 = get_half2_low(f0);
1156  __half a2 = get_half2_high(f0);
1157  __half b1 = get_half2_low(f1);
1158  __half b2 = get_half2_high(f1);
1159  f0 = combine_half(a1, b1);
1160  f1 = combine_half(a2, b2);
1161 }
1162 
1164 ptranspose(PacketBlock<Packet4h2,8>& kernel) {
1165  double* d_row0 = reinterpret_cast<double*>(&kernel.packet[0]);
1166  double* d_row1 = reinterpret_cast<double*>(&kernel.packet[1]);
1167  double* d_row2 = reinterpret_cast<double*>(&kernel.packet[2]);
1168  double* d_row3 = reinterpret_cast<double*>(&kernel.packet[3]);
1169  double* d_row4 = reinterpret_cast<double*>(&kernel.packet[4]);
1170  double* d_row5 = reinterpret_cast<double*>(&kernel.packet[5]);
1171  double* d_row6 = reinterpret_cast<double*>(&kernel.packet[6]);
1172  double* d_row7 = reinterpret_cast<double*>(&kernel.packet[7]);
1173  ptranspose_double(d_row0, d_row1, d_row2, d_row3,
1174  d_row4, d_row5, d_row6, d_row7);
1175 
1176 
1177  half2* f_row0 = reinterpret_cast<half2*>(d_row0);
1178  half2* f_row1 = reinterpret_cast<half2*>(d_row1);
1179  half2* f_row2 = reinterpret_cast<half2*>(d_row2);
1180  half2* f_row3 = reinterpret_cast<half2*>(d_row3);
1181  ptranspose_half2(f_row0, f_row1, f_row2, f_row3);
1182  ptranspose_half(f_row0[0], f_row1[0]);
1183  ptranspose_half(f_row0[1], f_row1[1]);
1184  ptranspose_half(f_row2[0], f_row3[0]);
1185  ptranspose_half(f_row2[1], f_row3[1]);
1186 
1187  f_row0 = reinterpret_cast<half2*>(d_row0 + 1);
1188  f_row1 = reinterpret_cast<half2*>(d_row1 + 1);
1189  f_row2 = reinterpret_cast<half2*>(d_row2 + 1);
1190  f_row3 = reinterpret_cast<half2*>(d_row3 + 1);
1191  ptranspose_half2(f_row0, f_row1, f_row2, f_row3);
1192  ptranspose_half(f_row0[0], f_row1[0]);
1193  ptranspose_half(f_row0[1], f_row1[1]);
1194  ptranspose_half(f_row2[0], f_row3[0]);
1195  ptranspose_half(f_row2[1], f_row3[1]);
1196 
1197  f_row0 = reinterpret_cast<half2*>(d_row4);
1198  f_row1 = reinterpret_cast<half2*>(d_row5);
1199  f_row2 = reinterpret_cast<half2*>(d_row6);
1200  f_row3 = reinterpret_cast<half2*>(d_row7);
1201  ptranspose_half2(f_row0, f_row1, f_row2, f_row3);
1202  ptranspose_half(f_row0[0], f_row1[0]);
1203  ptranspose_half(f_row0[1], f_row1[1]);
1204  ptranspose_half(f_row2[0], f_row3[0]);
1205  ptranspose_half(f_row2[1], f_row3[1]);
1206 
1207  f_row0 = reinterpret_cast<half2*>(d_row4 + 1);
1208  f_row1 = reinterpret_cast<half2*>(d_row5 + 1);
1209  f_row2 = reinterpret_cast<half2*>(d_row6 + 1);
1210  f_row3 = reinterpret_cast<half2*>(d_row7 + 1);
1211  ptranspose_half2(f_row0, f_row1, f_row2, f_row3);
1212  ptranspose_half(f_row0[0], f_row1[0]);
1213  ptranspose_half(f_row0[1], f_row1[1]);
1214  ptranspose_half(f_row2[0], f_row3[0]);
1215  ptranspose_half(f_row2[1], f_row3[1]);
1216 
1217 }
1218 
1219 template <>
1221 plset<Packet4h2>(const Eigen::half& a) {
1222 #if defined(EIGEN_HIP_DEVICE_COMPILE)
1223 
1224  Packet4h2 r;
1225  half2* p_alias = reinterpret_cast<half2*>(&r);
1226  p_alias[0] = __halves2half2(a, __hadd(a, __float2half(1.0f)));
1227  p_alias[1] = __halves2half2(__hadd(a, __float2half(2.0f)),
1228  __hadd(a, __float2half(3.0f)));
1229  p_alias[2] = __halves2half2(__hadd(a, __float2half(4.0f)),
1230  __hadd(a, __float2half(5.0f)));
1231  p_alias[3] = __halves2half2(__hadd(a, __float2half(6.0f)),
1232  __hadd(a, __float2half(7.0f)));
1233  return r;
1234 #elif defined(EIGEN_CUDA_HAS_FP16_ARITHMETIC)
1235  Packet4h2 r;
1236  half2* r_alias = reinterpret_cast<half2*>(&r);
1237 
1238  half2 b = pset1<half2>(a);
1239  half2 c;
1240  half2 half_offset0 = __halves2half2(__float2half(0.0f),__float2half(2.0f));
1241  half2 half_offset1 = __halves2half2(__float2half(4.0f),__float2half(6.0f));
1242 
1243  c = __hadd2(b, half_offset0);
1244  r_alias[0] = plset(__low2half(c));
1245  r_alias[1] = plset(__high2half(c));
1246 
1247  c = __hadd2(b, half_offset1);
1248  r_alias[2] = plset(__low2half(c));
1249  r_alias[3] = plset(__high2half(c));
1250 
1251  return r;
1252 
1253 #else
1254  float f = __half2float(a);
1255  Packet4h2 r;
1256  half2* p_alias = reinterpret_cast<half2*>(&r);
1257  p_alias[0] = combine_half(a, __float2half(f + 1.0f));
1258  p_alias[1] = combine_half(__float2half(f + 2.0f), __float2half(f + 3.0f));
1259  p_alias[2] = combine_half(__float2half(f + 4.0f), __float2half(f + 5.0f));
1260  p_alias[3] = combine_half(__float2half(f + 6.0f), __float2half(f + 7.0f));
1261  return r;
1262 #endif
1263 }
1264 
1265 template <>
1267 pselect<Packet4h2>(const Packet4h2& mask, const Packet4h2& a,
1268  const Packet4h2& b) {
1269  Packet4h2 r;
1270  half2* r_alias = reinterpret_cast<half2*>(&r);
1271  const half2* mask_alias = reinterpret_cast<const half2*>(&mask);
1272  const half2* a_alias = reinterpret_cast<const half2*>(&a);
1273  const half2* b_alias = reinterpret_cast<const half2*>(&b);
1274  r_alias[0] = pselect(mask_alias[0], a_alias[0], b_alias[0]);
1275  r_alias[1] = pselect(mask_alias[1], a_alias[1], b_alias[1]);
1276  r_alias[2] = pselect(mask_alias[2], a_alias[2], b_alias[2]);
1277  r_alias[3] = pselect(mask_alias[3], a_alias[3], b_alias[3]);
1278  return r;
1279 }
1280 
1281 template <>
1283 pcmp_eq<Packet4h2>(const Packet4h2& a, const Packet4h2& b) {
1284  Packet4h2 r;
1285  half2* r_alias = reinterpret_cast<half2*>(&r);
1286  const half2* a_alias = reinterpret_cast<const half2*>(&a);
1287  const half2* b_alias = reinterpret_cast<const half2*>(&b);
1288  r_alias[0] = pcmp_eq(a_alias[0], b_alias[0]);
1289  r_alias[1] = pcmp_eq(a_alias[1], b_alias[1]);
1290  r_alias[2] = pcmp_eq(a_alias[2], b_alias[2]);
1291  r_alias[3] = pcmp_eq(a_alias[3], b_alias[3]);
1292  return r;
1293 }
1294 
1295 template <>
1296 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 pand<Packet4h2>(
1297  const Packet4h2& a, const Packet4h2& b) {
1298  Packet4h2 r;
1299  half2* r_alias = reinterpret_cast<half2*>(&r);
1300  const half2* a_alias = reinterpret_cast<const half2*>(&a);
1301  const half2* b_alias = reinterpret_cast<const half2*>(&b);
1302  r_alias[0] = pand(a_alias[0], b_alias[0]);
1303  r_alias[1] = pand(a_alias[1], b_alias[1]);
1304  r_alias[2] = pand(a_alias[2], b_alias[2]);
1305  r_alias[3] = pand(a_alias[3], b_alias[3]);
1306  return r;
1307 }
1308 
1309 template <>
1310 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 por<Packet4h2>(
1311  const Packet4h2& a, const Packet4h2& b) {
1312  Packet4h2 r;
1313  half2* r_alias = reinterpret_cast<half2*>(&r);
1314  const half2* a_alias = reinterpret_cast<const half2*>(&a);
1315  const half2* b_alias = reinterpret_cast<const half2*>(&b);
1316  r_alias[0] = por(a_alias[0], b_alias[0]);
1317  r_alias[1] = por(a_alias[1], b_alias[1]);
1318  r_alias[2] = por(a_alias[2], b_alias[2]);
1319  r_alias[3] = por(a_alias[3], b_alias[3]);
1320  return r;
1321 }
1322 
1323 template <>
1324 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 pxor<Packet4h2>(
1325  const Packet4h2& a, const Packet4h2& b) {
1326  Packet4h2 r;
1327  half2* r_alias = reinterpret_cast<half2*>(&r);
1328  const half2* a_alias = reinterpret_cast<const half2*>(&a);
1329  const half2* b_alias = reinterpret_cast<const half2*>(&b);
1330  r_alias[0] = pxor(a_alias[0], b_alias[0]);
1331  r_alias[1] = pxor(a_alias[1], b_alias[1]);
1332  r_alias[2] = pxor(a_alias[2], b_alias[2]);
1333  r_alias[3] = pxor(a_alias[3], b_alias[3]);
1334  return r;
1335 }
1336 
1337 template <>
1339 pandnot<Packet4h2>(const Packet4h2& a, const Packet4h2& b) {
1340  Packet4h2 r;
1341  half2* r_alias = reinterpret_cast<half2*>(&r);
1342  const half2* a_alias = reinterpret_cast<const half2*>(&a);
1343  const half2* b_alias = reinterpret_cast<const half2*>(&b);
1344  r_alias[0] = pandnot(a_alias[0], b_alias[0]);
1345  r_alias[1] = pandnot(a_alias[1], b_alias[1]);
1346  r_alias[2] = pandnot(a_alias[2], b_alias[2]);
1347  r_alias[3] = pandnot(a_alias[3], b_alias[3]);
1348  return r;
1349 }
1350 
1351 template <>
1352 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 padd<Packet4h2>(
1353  const Packet4h2& a, const Packet4h2& b) {
1354  Packet4h2 r;
1355  half2* r_alias = reinterpret_cast<half2*>(&r);
1356  const half2* a_alias = reinterpret_cast<const half2*>(&a);
1357  const half2* b_alias = reinterpret_cast<const half2*>(&b);
1358  r_alias[0] = padd(a_alias[0], b_alias[0]);
1359  r_alias[1] = padd(a_alias[1], b_alias[1]);
1360  r_alias[2] = padd(a_alias[2], b_alias[2]);
1361  r_alias[3] = padd(a_alias[3], b_alias[3]);
1362  return r;
1363 }
1364 
1365 template <>
1366 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 psub<Packet4h2>(
1367  const Packet4h2& a, const Packet4h2& b) {
1368  Packet4h2 r;
1369  half2* r_alias = reinterpret_cast<half2*>(&r);
1370  const half2* a_alias = reinterpret_cast<const half2*>(&a);
1371  const half2* b_alias = reinterpret_cast<const half2*>(&b);
1372  r_alias[0] = psub(a_alias[0], b_alias[0]);
1373  r_alias[1] = psub(a_alias[1], b_alias[1]);
1374  r_alias[2] = psub(a_alias[2], b_alias[2]);
1375  r_alias[3] = psub(a_alias[3], b_alias[3]);
1376  return r;
1377 }
1378 
1379 template <>
1380 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 pnegate(const Packet4h2& a) {
1381  Packet4h2 r;
1382  half2* r_alias = reinterpret_cast<half2*>(&r);
1383  const half2* a_alias = reinterpret_cast<const half2*>(&a);
1384  r_alias[0] = pnegate(a_alias[0]);
1385  r_alias[1] = pnegate(a_alias[1]);
1386  r_alias[2] = pnegate(a_alias[2]);
1387  r_alias[3] = pnegate(a_alias[3]);
1388  return r;
1389 }
1390 
1391 template <>
1392 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 pconj(const Packet4h2& a) {
1393  return a;
1394 }
1395 
1396 template <>
1397 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 pmul<Packet4h2>(
1398  const Packet4h2& a, const Packet4h2& b) {
1399  Packet4h2 r;
1400  half2* r_alias = reinterpret_cast<half2*>(&r);
1401  const half2* a_alias = reinterpret_cast<const half2*>(&a);
1402  const half2* b_alias = reinterpret_cast<const half2*>(&b);
1403  r_alias[0] = pmul(a_alias[0], b_alias[0]);
1404  r_alias[1] = pmul(a_alias[1], b_alias[1]);
1405  r_alias[2] = pmul(a_alias[2], b_alias[2]);
1406  r_alias[3] = pmul(a_alias[3], b_alias[3]);
1407  return r;
1408 }
1409 
1410 template <>
1411 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 pmadd<Packet4h2>(
1412  const Packet4h2& a, const Packet4h2& b, const Packet4h2& c) {
1413  Packet4h2 r;
1414  half2* r_alias = reinterpret_cast<half2*>(&r);
1415  const half2* a_alias = reinterpret_cast<const half2*>(&a);
1416  const half2* b_alias = reinterpret_cast<const half2*>(&b);
1417  const half2* c_alias = reinterpret_cast<const half2*>(&c);
1418  r_alias[0] = pmadd(a_alias[0], b_alias[0], c_alias[0]);
1419  r_alias[1] = pmadd(a_alias[1], b_alias[1], c_alias[1]);
1420  r_alias[2] = pmadd(a_alias[2], b_alias[2], c_alias[2]);
1421  r_alias[3] = pmadd(a_alias[3], b_alias[3], c_alias[3]);
1422  return r;
1423 }
1424 
1425 template <>
1426 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 pdiv<Packet4h2>(
1427  const Packet4h2& a, const Packet4h2& b) {
1428  Packet4h2 r;
1429  half2* r_alias = reinterpret_cast<half2*>(&r);
1430  const half2* a_alias = reinterpret_cast<const half2*>(&a);
1431  const half2* b_alias = reinterpret_cast<const half2*>(&b);
1432  r_alias[0] = pdiv(a_alias[0], b_alias[0]);
1433  r_alias[1] = pdiv(a_alias[1], b_alias[1]);
1434  r_alias[2] = pdiv(a_alias[2], b_alias[2]);
1435  r_alias[3] = pdiv(a_alias[3], b_alias[3]);
1436  return r;
1437 }
1438 
1439 template <>
1440 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 pmin<Packet4h2>(
1441  const Packet4h2& a, const Packet4h2& b) {
1442  Packet4h2 r;
1443  half2* r_alias = reinterpret_cast<half2*>(&r);
1444  const half2* a_alias = reinterpret_cast<const half2*>(&a);
1445  const half2* b_alias = reinterpret_cast<const half2*>(&b);
1446  r_alias[0] = pmin(a_alias[0], b_alias[0]);
1447  r_alias[1] = pmin(a_alias[1], b_alias[1]);
1448  r_alias[2] = pmin(a_alias[2], b_alias[2]);
1449  r_alias[3] = pmin(a_alias[3], b_alias[3]);
1450  return r;
1451 }
1452 
1453 template <>
1454 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 pmax<Packet4h2>(
1455  const Packet4h2& a, const Packet4h2& b) {
1456  Packet4h2 r;
1457  half2* r_alias = reinterpret_cast<half2*>(&r);
1458  const half2* a_alias = reinterpret_cast<const half2*>(&a);
1459  const half2* b_alias = reinterpret_cast<const half2*>(&b);
1460  r_alias[0] = pmax(a_alias[0], b_alias[0]);
1461  r_alias[1] = pmax(a_alias[1], b_alias[1]);
1462  r_alias[2] = pmax(a_alias[2], b_alias[2]);
1463  r_alias[3] = pmax(a_alias[3], b_alias[3]);
1464  return r;
1465 }
1466 
1467 template <>
1469  const Packet4h2& a) {
1470  const half2* a_alias = reinterpret_cast<const half2*>(&a);
1471 
1472  return predux(a_alias[0]) + predux(a_alias[1]) +
1473  predux(a_alias[2]) + predux(a_alias[3]);
1474 }
1475 
1476 template <>
1477 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::half predux_max<Packet4h2>(
1478  const Packet4h2& a) {
1479  const half2* a_alias = reinterpret_cast<const half2*>(&a);
1480  half2 m0 = combine_half(predux_max(a_alias[0]),
1481  predux_max(a_alias[1]));
1482  half2 m1 = combine_half(predux_max(a_alias[2]),
1483  predux_max(a_alias[3]));
1484  __half first = predux_max(m0);
1485  __half second = predux_max(m1);
1486 #if defined(EIGEN_CUDA_HAS_FP16_ARITHMETIC)
1487  return (__hgt(first, second) ? first : second);
1488 #else
1489  float ffirst = __half2float(first);
1490  float fsecond = __half2float(second);
1491  return (ffirst > fsecond)? first: second;
1492 #endif
1493 }
1494 
1495 template <>
1496 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::half predux_min<Packet4h2>(
1497  const Packet4h2& a) {
1498  const half2* a_alias = reinterpret_cast<const half2*>(&a);
1499  half2 m0 = combine_half(predux_min(a_alias[0]),
1500  predux_min(a_alias[1]));
1501  half2 m1 = combine_half(predux_min(a_alias[2]),
1502  predux_min(a_alias[3]));
1503  __half first = predux_min(m0);
1504  __half second = predux_min(m1);
1505 #if defined(EIGEN_CUDA_HAS_FP16_ARITHMETIC)
1506  return (__hlt(first, second) ? first : second);
1507 #else
1508  float ffirst = __half2float(first);
1509  float fsecond = __half2float(second);
1510  return (ffirst < fsecond)? first: second;
1511 #endif
1512 }
1513 
1514 // likely overflow/underflow
1515 template <>
1516 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::half predux_mul<Packet4h2>(
1517  const Packet4h2& a) {
1518  const half2* a_alias = reinterpret_cast<const half2*>(&a);
1519  return predux_mul(pmul(pmul(a_alias[0], a_alias[1]),
1520  pmul(a_alias[2], a_alias[3])));
1521 }
1522 
1523 template <>
1525 plog1p<Packet4h2>(const Packet4h2& a) {
1526  Packet4h2 r;
1527  half2* r_alias = reinterpret_cast<half2*>(&r);
1528  const half2* a_alias = reinterpret_cast<const half2*>(&a);
1529  r_alias[0] = plog1p(a_alias[0]);
1530  r_alias[1] = plog1p(a_alias[1]);
1531  r_alias[2] = plog1p(a_alias[2]);
1532  r_alias[3] = plog1p(a_alias[3]);
1533  return r;
1534 }
1535 
1536 template <>
1538 pexpm1<Packet4h2>(const Packet4h2& a) {
1539  Packet4h2 r;
1540  half2* r_alias = reinterpret_cast<half2*>(&r);
1541  const half2* a_alias = reinterpret_cast<const half2*>(&a);
1542  r_alias[0] = pexpm1(a_alias[0]);
1543  r_alias[1] = pexpm1(a_alias[1]);
1544  r_alias[2] = pexpm1(a_alias[2]);
1545  r_alias[3] = pexpm1(a_alias[3]);
1546  return r;
1547 }
1548 
1549 template <>
1550 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 plog<Packet4h2>(const Packet4h2& a) {
1551  Packet4h2 r;
1552  half2* r_alias = reinterpret_cast<half2*>(&r);
1553  const half2* a_alias = reinterpret_cast<const half2*>(&a);
1554  r_alias[0] = plog(a_alias[0]);
1555  r_alias[1] = plog(a_alias[1]);
1556  r_alias[2] = plog(a_alias[2]);
1557  r_alias[3] = plog(a_alias[3]);
1558  return r;
1559 }
1560 
1561 template <>
1562 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 pexp<Packet4h2>(const Packet4h2& a) {
1563  Packet4h2 r;
1564  half2* r_alias = reinterpret_cast<half2*>(&r);
1565  const half2* a_alias = reinterpret_cast<const half2*>(&a);
1566  r_alias[0] = pexp(a_alias[0]);
1567  r_alias[1] = pexp(a_alias[1]);
1568  r_alias[2] = pexp(a_alias[2]);
1569  r_alias[3] = pexp(a_alias[3]);
1570  return r;
1571 }
1572 
1573 template <>
1574 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 psqrt<Packet4h2>(const Packet4h2& a) {
1575  Packet4h2 r;
1576  half2* r_alias = reinterpret_cast<half2*>(&r);
1577  const half2* a_alias = reinterpret_cast<const half2*>(&a);
1578  r_alias[0] = psqrt(a_alias[0]);
1579  r_alias[1] = psqrt(a_alias[1]);
1580  r_alias[2] = psqrt(a_alias[2]);
1581  r_alias[3] = psqrt(a_alias[3]);
1582  return r;
1583 }
1584 
1585 template <>
1587 prsqrt<Packet4h2>(const Packet4h2& a) {
1588  Packet4h2 r;
1589  half2* r_alias = reinterpret_cast<half2*>(&r);
1590  const half2* a_alias = reinterpret_cast<const half2*>(&a);
1591  r_alias[0] = prsqrt(a_alias[0]);
1592  r_alias[1] = prsqrt(a_alias[1]);
1593  r_alias[2] = prsqrt(a_alias[2]);
1594  r_alias[3] = prsqrt(a_alias[3]);
1595  return r;
1596 }
1597 
1598 // The following specialized padd, pmul, pdiv, pmin, pmax, pset1 are needed for
1599 // the implementation of GPU half reduction.
1600 template<>
1601 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 padd<half2>(const half2& a,
1602  const half2& b) {
1603 #if defined(EIGEN_GPU_HAS_FP16_ARITHMETIC)
1604  return __hadd2(a, b);
1605 #else
1606  float a1 = __low2float(a);
1607  float a2 = __high2float(a);
1608  float b1 = __low2float(b);
1609  float b2 = __high2float(b);
1610  float r1 = a1 + b1;
1611  float r2 = a2 + b2;
1612  return __floats2half2_rn(r1, r2);
1613 #endif
1614 }
1615 
1616 template<>
1617 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pmul<half2>(const half2& a,
1618  const half2& b) {
1619 #if defined(EIGEN_GPU_HAS_FP16_ARITHMETIC)
1620  return __hmul2(a, b);
1621 #else
1622  float a1 = __low2float(a);
1623  float a2 = __high2float(a);
1624  float b1 = __low2float(b);
1625  float b2 = __high2float(b);
1626  float r1 = a1 * b1;
1627  float r2 = a2 * b2;
1628  return __floats2half2_rn(r1, r2);
1629 #endif
1630 }
1631 
1632 template<>
1633 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pdiv<half2>(const half2& a,
1634  const half2& b) {
1635 #if defined(EIGEN_GPU_HAS_FP16_ARITHMETIC)
1636  return __h2div(a, b);
1637 #else
1638  float a1 = __low2float(a);
1639  float a2 = __high2float(a);
1640  float b1 = __low2float(b);
1641  float b2 = __high2float(b);
1642  float r1 = a1 / b1;
1643  float r2 = a2 / b2;
1644  return __floats2half2_rn(r1, r2);
1645 #endif
1646 }
1647 
1648 template<>
1649 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pmin<half2>(const half2& a,
1650  const half2& b) {
1651  float a1 = __low2float(a);
1652  float a2 = __high2float(a);
1653  float b1 = __low2float(b);
1654  float b2 = __high2float(b);
1655  __half r1 = a1 < b1 ? get_half2_low(a) : get_half2_low(b);
1656  __half r2 = a2 < b2 ? get_half2_high(a) : get_half2_high(b);
1657  return combine_half(r1, r2);
1658 }
1659 
1660 template<>
1661 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pmax<half2>(const half2& a,
1662  const half2& b) {
1663  float a1 = __low2float(a);
1664  float a2 = __high2float(a);
1665  float b1 = __low2float(b);
1666  float b2 = __high2float(b);
1667  __half r1 = a1 > b1 ? get_half2_low(a) : get_half2_low(b);
1668  __half r2 = a2 > b2 ? get_half2_high(a) : get_half2_high(b);
1669  return combine_half(r1, r2);
1670 }
1671 
1672 // #endif // defined(EIGEN_CUDA_ARCH) || defined(EIGEN_HIPCC) || (defined(EIGEN_CUDACC) && EIGEN_COMP_CLANG && !EIGEN_COMP_NVCC)
1673 
1674 #endif // defined(EIGEN_HAS_CUDA_FP16) || defined(EIGEN_HAS_HIP_FP16)
1675 
1676 #undef EIGEN_GPU_HAS_LDG
1677 #undef EIGEN_CUDA_HAS_FP16_ARITHMETIC
1678 #undef EIGEN_GPU_HAS_FP16_ARITHMETIC
1679 
1680 } // end namespace internal
1681 
1682 } // end namespace Eigen
1683 
1684 
1685 #endif // EIGEN_PACKET_MATH_GPU_H
Eigen::internal::pcmp_eq
EIGEN_STRONG_INLINE Packet2cf pcmp_eq(const Packet2cf &a, const Packet2cf &b)
Definition: AltiVec/Complex.h:231
EIGEN_DEVICE_FUNC
#define EIGEN_DEVICE_FUNC
Definition: Macros.h:976
Eigen
Namespace containing all symbols from the Eigen library.
Definition: jet.h:637
Eigen::internal::default_packet_traits::HasIGammac
@ HasIGammac
Definition: GenericPacketMath.h:95
Eigen::internal::psqrt
EIGEN_STRONG_INLINE Packet4f psqrt(const Packet4f &a)
Definition: MSA/PacketMath.h:723
simple_graph::b1
Vector2 b1(2, -1)
Eigen::internal::default_packet_traits::HasBetaInc
@ HasBetaInc
Definition: GenericPacketMath.h:96
test_constructor::f1
auto f1
Definition: testHybridNonlinearFactor.cpp:56
Eigen::internal::default_packet_traits::HasLog
@ HasLog
Definition: GenericPacketMath.h:70
Eigen::internal::packet_traits::HasSub
@ HasSub
Definition: GenericPacketMath.h:118
double2
Definition: dd_real.h:74
Eigen::internal::pstoreu
EIGEN_DEVICE_FUNC void pstoreu(Scalar *to, const Packet &from)
Definition: GenericPacketMath.h:700
Eigen::internal::packet_traits::size
@ size
Definition: GenericPacketMath.h:112
r2
static const double r2
Definition: testSmartRangeFactor.cpp:32
c
Scalar Scalar * c
Definition: benchVecAdd.cpp:17
b
Scalar * b
Definition: benchVecAdd.cpp:17
Eigen::internal::pcmp_lt
EIGEN_STRONG_INLINE Packet4f pcmp_lt(const Packet4f &a, const Packet4f &b)
Definition: AltiVec/PacketMath.h:868
Eigen::internal::pfirst
EIGEN_STRONG_INLINE bfloat16 pfirst(const Packet8bf &a)
Definition: AltiVec/PacketMath.h:1429
simple_graph::b2
Vector2 b2(4, -5)
Eigen::internal::default_packet_traits::HasLog1p
@ HasLog1p
Definition: GenericPacketMath.h:71
m1
Matrix3d m1
Definition: IOFormat.cpp:2
Eigen::internal::default_packet_traits::HasIGammaDerA
@ HasIGammaDerA
Definition: GenericPacketMath.h:93
Eigen::internal::packet_traits::type
T type
Definition: GenericPacketMath.h:108
Eigen::internal::predux_mul
EIGEN_DEVICE_FUNC unpacket_traits< Packet >::type predux_mul(const Packet &a)
Definition: GenericPacketMath.h:882
Eigen::internal::pdiv
EIGEN_DEVICE_FUNC Packet pdiv(const Packet &a, const Packet &b)
Definition: GenericPacketMath.h:244
m0
static const DiscreteKey m0(M(0), 2)
Eigen::internal::pconj
EIGEN_STRONG_INLINE Packet2cf pconj(const Packet2cf &a)
Definition: AltiVec/Complex.h:167
Eigen::internal::default_packet_traits::HasPolygamma
@ HasPolygamma
Definition: GenericPacketMath.h:87
Eigen::internal::unpacket_traits::vectorizable
@ vectorizable
Definition: GenericPacketMath.h:140
Eigen::internal::default_packet_traits::HasBlend
@ HasBlend
Definition: GenericPacketMath.h:60
Eigen::internal::default_packet_traits::HasCos
@ HasCos
Definition: GenericPacketMath.h:76
r1
static const double r1
Definition: testSmartRangeFactor.cpp:32
Eigen::internal::default_packet_traits::HasDiGamma
@ HasDiGamma
Definition: GenericPacketMath.h:85
Eigen::internal::default_packet_traits::HasZeta
@ HasZeta
Definition: GenericPacketMath.h:86
Eigen::internal::pxor
EIGEN_STRONG_INLINE Packet8h pxor(const Packet8h &a, const Packet8h &b)
Definition: AVX/PacketMath.h:1047
Eigen::internal::packet_traits::HasHalfPacket
@ HasHalfPacket
Definition: GenericPacketMath.h:114
Eigen::internal::is_arithmetic::value
@ value
Definition: Meta.h:133
Eigen::internal::unpacket_traits::masked_store_available
@ masked_store_available
Definition: GenericPacketMath.h:142
boost::multiprecision::fabs
Real fabs(const Real &a)
Definition: boostmultiprec.cpp:119
c1
static double c1
Definition: airy.c:54
double2::x
double x[2]
Definition: dd_real.h:76
align_3::a1
Point2 a1
Definition: testPose2.cpp:769
Eigen::internal::pgather
EIGEN_DEVICE_FUNC Packet pgather(const Scalar *from, Index)
Definition: GenericPacketMath.h:712
Eigen::internal::unpacket_traits::size
@ size
Definition: GenericPacketMath.h:138
Eigen::internal::unpacket_traits::alignment
@ alignment
Definition: GenericPacketMath.h:139
Eigen::internal::pstore< double >
EIGEN_STRONG_INLINE void pstore< double >(double *to, const Packet4d &from)
Definition: AVX/PacketMath.h:623
Eigen::internal::unpacket_traits::half
T half
Definition: GenericPacketMath.h:135
Eigen::internal::default_packet_traits::HasGammaSampleDerAlpha
@ HasGammaSampleDerAlpha
Definition: GenericPacketMath.h:94
test_constructor::f0
auto f0
Definition: testHybridNonlinearFactor.cpp:55
Eigen::internal::default_packet_traits::HasSin
@ HasSin
Definition: GenericPacketMath.h:75
Eigen::internal::packet_traits::AlignedOnScalar
@ AlignedOnScalar
Definition: GenericPacketMath.h:113
Eigen::internal::default_packet_traits::HasNdtri
@ HasNdtri
Definition: GenericPacketMath.h:90
Eigen::internal::packet_traits::half
T half
Definition: GenericPacketMath.h:109
Eigen::internal::pnegate
EIGEN_STRONG_INLINE Packet2cf pnegate(const Packet2cf &a)
Definition: AltiVec/Complex.h:166
Eigen::internal::default_packet_traits::HasBessel
@ HasBessel
Definition: GenericPacketMath.h:91
Eigen::internal::ptranspose
EIGEN_STRONG_INLINE void ptranspose(PacketBlock< Packet2cf, 2 > &kernel)
Definition: AltiVec/Complex.h:224
Eigen::internal::pselect
EIGEN_STRONG_INLINE Packet4f pselect(const Packet4f &mask, const Packet4f &a, const Packet4f &b)
Definition: AltiVec/PacketMath.h:917
Eigen::internal::pexpm1
EIGEN_DECLARE_FUNCTION_ALLOWING_MULTIPLE_DEFINITIONS Packet pexpm1(const Packet &a)
Definition: GenericPacketMath.h:792
EIGEN_STRONG_INLINE
#define EIGEN_STRONG_INLINE
Definition: Macros.h:917
Eigen::internal::first
EIGEN_CONSTEXPR Index first(const T &x) EIGEN_NOEXCEPT
Definition: IndexedViewHelper.h:81
Eigen::internal::pmax
EIGEN_DEVICE_FUNC Packet pmax(const Packet &a, const Packet &b)
Definition: GenericPacketMath.h:524
Eigen::internal::default_packet_traits::HasErf
@ HasErf
Definition: GenericPacketMath.h:88
Eigen::internal::pscatter
EIGEN_DEVICE_FUNC void pscatter(Scalar *to, const Packet &from, Index)
Definition: GenericPacketMath.h:715
Eigen::internal::default_packet_traits::HasSqrt
@ HasSqrt
Definition: GenericPacketMath.h:66
Eigen::internal::pand
EIGEN_STRONG_INLINE Packet8h pand(const Packet8h &a, const Packet8h &b)
Definition: AVX/PacketMath.h:1050
Eigen::internal::psub
EIGEN_DEVICE_FUNC Packet psub(const Packet &a, const Packet &b)
Definition: GenericPacketMath.h:222
EIGEN_ALWAYS_INLINE
#define EIGEN_ALWAYS_INLINE
Definition: Macros.h:932
Eigen::internal::default_packet_traits::HasFloor
@ HasFloor
Definition: GenericPacketMath.h:100
Eigen::internal::ploaddup
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet ploaddup(const typename unpacket_traits< Packet >::type *from)
Definition: GenericPacketMath.h:631
Eigen::bfloat16_impl::fmin
EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 fmin(const bfloat16 &a, const bfloat16 &b)
Definition: BFloat16.h:582
Eigen::internal::pexp
EIGEN_DECLARE_FUNCTION_ALLOWING_MULTIPLE_DEFINITIONS Packet pexp(const Packet &a)
Definition: GenericPacketMath.h:788
Eigen::internal::pstore
EIGEN_DEVICE_FUNC void pstore(Scalar *to, const Packet &from)
Definition: GenericPacketMath.h:696
Eigen::internal::unpacket_traits::masked_load_available
@ masked_load_available
Definition: GenericPacketMath.h:141
tree::f
Point2(* f)(const Point3 &, OptionalJacobian< 2, 3 >)
Definition: testExpression.cpp:218
Eigen::internal::pzero
EIGEN_STRONG_INLINE Packet8f pzero(const Packet8f &)
Definition: AVX/PacketMath.h:247
Eigen::Aligned16
@ Aligned16
Definition: Constants.h:235
a
ArrayXXi a
Definition: Array_initializer_list_23_cxx11.cpp:1
Eigen::internal::default_packet_traits::HasExp
@ HasExp
Definition: GenericPacketMath.h:68
Eigen::internal::predux
EIGEN_DEVICE_FUNC unpacket_traits< Packet >::type predux(const Packet &a)
Definition: GenericPacketMath.h:875
Eigen::internal::pmul
EIGEN_DEVICE_FUNC Packet pmul(const Packet &a, const Packet &b)
Definition: GenericPacketMath.h:237
Eigen::internal::unpacket_traits::type
T type
Definition: GenericPacketMath.h:134
Eigen::internal::default_packet_traits::HasRsqrt
@ HasRsqrt
Definition: GenericPacketMath.h:67
Eigen::internal::default_packet_traits::HasIGamma
@ HasIGamma
Definition: GenericPacketMath.h:92
align_3::a2
Point2 a2
Definition: testPose2.cpp:770
Eigen::internal::default_packet_traits::HasExpm1
@ HasExpm1
Definition: GenericPacketMath.h:69
Eigen::bfloat16_impl::fmax
EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bfloat16 fmax(const bfloat16 &a, const bfloat16 &b)
Definition: BFloat16.h:587
Eigen::internal::pmadd
EIGEN_STRONG_INLINE Packet4f pmadd(const Packet4f &a, const Packet4f &b, const Packet4f &c)
Definition: AltiVec/PacketMath.h:827
Eigen::internal::plog
EIGEN_DECLARE_FUNCTION_ALLOWING_MULTIPLE_DEFINITIONS Packet plog(const Packet &a)
Definition: GenericPacketMath.h:796
Eigen::internal::pstoreu< double >
EIGEN_STRONG_INLINE void pstoreu< double >(double *to, const Packet4d &from)
Definition: AVX/PacketMath.h:627
Eigen::half_impl::raw_uint16_to_half
EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR __half_raw raw_uint16_to_half(numext::uint16_t x)
Definition: Half.h:495
Eigen::internal::packet_traits::HasAdd
@ HasAdd
Definition: GenericPacketMath.h:117
c2
static double c2
Definition: airy.c:55
Eigen::internal::pabs
EIGEN_STRONG_INLINE Packet4f pabs(const Packet4f &a)
Definition: AltiVec/PacketMath.h:1176
Eigen::internal::padd
EIGEN_DEVICE_FUNC Packet padd(const Packet &a, const Packet &b)
Definition: GenericPacketMath.h:215
Eigen::internal::pmin
EIGEN_DEVICE_FUNC Packet pmin(const Packet &a, const Packet &b)
Definition: GenericPacketMath.h:512
Eigen::internal::plset
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet plset(const typename unpacket_traits< Packet >::type &a)
Definition: GenericPacketMath.h:679
Eigen::TensorSycl::internal::scan_step::second
@ second
gtsam.examples.DogLegOptimizerExample.float
float
Definition: DogLegOptimizerExample.py:113
Eigen::internal::por
EIGEN_STRONG_INLINE Packet8h por(const Packet8h &a, const Packet8h &b)
Definition: AVX/PacketMath.h:1042
Eigen::internal::default_packet_traits::HasErfc
@ HasErfc
Definition: GenericPacketMath.h:89
Eigen::internal::default_packet_traits::HasDiv
@ HasDiv
Definition: GenericPacketMath.h:65
Eigen::internal::plog1p
EIGEN_DECLARE_FUNCTION_ALLOWING_MULTIPLE_DEFINITIONS Packet plog1p(const Packet &a)
Definition: GenericPacketMath.h:800
internal
Definition: BandTriangularSolver.h:13
Eigen::internal::pstoreu< float >
EIGEN_STRONG_INLINE void pstoreu< float >(float *to, const Packet4f &from)
Definition: AltiVec/PacketMath.h:1088
Vector2::x
float x
Definition: test_operator_overloading.cpp:94
Eigen::internal::prsqrt
EIGEN_STRONG_INLINE Packet4f prsqrt(const Packet4f &a)
Definition: MSA/PacketMath.h:730
Eigen::half
Definition: Half.h:142
Eigen::internal::pandnot
EIGEN_STRONG_INLINE Packet8h pandnot(const Packet8h &a, const Packet8h &b)
Definition: AVX/PacketMath.h:1053
Eigen::internal::predux_min
EIGEN_DEVICE_FUNC unpacket_traits< Packet >::type predux_min(const Packet &a)
Definition: GenericPacketMath.h:890
Eigen::internal::packet_traits::HasMul
@ HasMul
Definition: GenericPacketMath.h:119
Eigen::internal::packet_traits::Vectorizable
@ Vectorizable
Definition: GenericPacketMath.h:111
Eigen::internal::ploadu
EIGEN_DEVICE_FUNC Packet ploadu(const typename unpacket_traits< Packet >::type *from)
Definition: GenericPacketMath.h:603
Eigen::internal::pload
EIGEN_DEVICE_FUNC Packet pload(const typename unpacket_traits< Packet >::type *from)
Definition: GenericPacketMath.h:599
Eigen::internal::pstore< float >
EIGEN_STRONG_INLINE void pstore< float >(float *to, const Packet4f &from)
Definition: AltiVec/PacketMath.h:491
Eigen::internal::ptrue
EIGEN_STRONG_INLINE Packet8h ptrue(const Packet8h &a)
Definition: AVX/PacketMath.h:978
floor
const EIGEN_DEVICE_FUNC FloorReturnType floor() const
Definition: ArrayCwiseUnaryOps.h:481
Eigen::internal::predux_max
EIGEN_DEVICE_FUNC unpacket_traits< Packet >::type predux_max(const Packet &a)
Definition: GenericPacketMath.h:905
Eigen::Index
EIGEN_DEFAULT_DENSE_INDEX_TYPE Index
The Index type as used for the API.
Definition: Meta.h:74
Eigen::internal::default_packet_traits::HasLGamma
@ HasLGamma
Definition: GenericPacketMath.h:84


gtsam
Author(s):
autogenerated on Fri Nov 1 2024 03:33:53