Eigen  3.2.93
PacketMathHalf.h
1 // This file is part of Eigen, a lightweight C++ template library
2 // for linear algebra.
3 //
4 // Copyright (C) 2016 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_HALF_CUDA_H
11 #define EIGEN_PACKET_MATH_HALF_CUDA_H
12 
13 
14 namespace Eigen {
15 namespace internal {
16 
17 // Most of the following operations require arch >= 3.0
18 #if defined(EIGEN_HAS_CUDA_FP16) && defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 300
19 
20 template<> struct is_arithmetic<half2> { enum { value = true }; };
21 
22 template<> struct packet_traits<Eigen::half> : default_packet_traits
23 {
24  typedef half2 type;
25  typedef half2 half;
26  enum {
27  Vectorizable = 1,
28  AlignedOnScalar = 1,
29  size=2,
30  HasHalfPacket = 0,
31  HasAdd = 1,
32  HasMul = 1,
33  HasDiv = 1,
34  HasSqrt = 1,
35  HasRsqrt = 1,
36  HasExp = 1,
37  HasLog = 1
38  };
39 };
40 
41 template<> struct unpacket_traits<half2> { typedef Eigen::half type; enum {size=2, alignment=Aligned16}; typedef half2 half; };
42 
43 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pset1<half2>(const Eigen::half& from) {
44  return __half2half2(from);
45 }
46 
47 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pload<half2>(const Eigen::half* from) {
48  return *reinterpret_cast<const half2*>(from);
49 }
50 
51 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 ploadu<half2>(const Eigen::half* from) {
52  return __halves2half2(from[0], from[1]);
53 }
54 
55 template<> EIGEN_STRONG_INLINE half2 ploaddup<half2>(const Eigen::half* from) {
56  return __halves2half2(from[0], from[0]);
57 }
58 
59 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void pstore<Eigen::half>(Eigen::half* to, const half2& from) {
60  *reinterpret_cast<half2*>(to) = from;
61 }
62 
63 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void pstoreu<Eigen::half>(Eigen::half* to, const half2& from) {
64  to[0] = __low2half(from);
65  to[1] = __high2half(from);
66 }
67 
68 template<>
69  EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE half2 ploadt_ro<half2, Aligned>(const Eigen::half* from) {
70 #if __CUDA_ARCH__ >= 350
71  return __ldg((const half2*)from);
72 #else
73  return __halves2half2(*(from+0), *(from+1));
74 #endif
75 }
76 
77 template<>
78 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE half2 ploadt_ro<half2, Unaligned>(const Eigen::half* from) {
79 #if __CUDA_ARCH__ >= 350
80  return __halves2half2(__ldg(from+0), __ldg(from+1));
81 #else
82  return __halves2half2(*(from+0), *(from+1));
83 #endif
84 }
85 
86 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pgather<Eigen::half, half2>(const Eigen::half* from, Index stride) {
87  return __halves2half2(from[0*stride], from[1*stride]);
88 }
89 
90 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void pscatter<Eigen::half, half2>(Eigen::half* to, const half2& from, Index stride) {
91  to[stride*0] = __low2half(from);
92  to[stride*1] = __high2half(from);
93 }
94 
95 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::half pfirst<half2>(const half2& a) {
96  return __low2half(a);
97 }
98 
99 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pabs<half2>(const half2& a) {
100  half2 result;
101  result.x = a.x & 0x7FFF7FFF;
102  return result;
103 }
104 
105 
106 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void
107 ptranspose(PacketBlock<half2,2>& kernel) {
108  __half a1 = __low2half(kernel.packet[0]);
109  __half a2 = __high2half(kernel.packet[0]);
110  __half b1 = __low2half(kernel.packet[1]);
111  __half b2 = __high2half(kernel.packet[1]);
112  kernel.packet[0] = __halves2half2(a1, b1);
113  kernel.packet[1] = __halves2half2(a2, b2);
114 }
115 
116 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 plset<half2>(const Eigen::half& a) {
117 #if __CUDA_ARCH__ >= 530
118  return __halves2half2(a, __hadd(a, __float2half(1.0f)));
119 #else
120  float f = __half2float(a) + 1.0f;
121  return __halves2half2(a, __float2half(f));
122 #endif
123 }
124 
125 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 padd<half2>(const half2& a, const half2& b) {
126 #if __CUDA_ARCH__ >= 530
127  return __hadd2(a, b);
128 #else
129  float a1 = __low2float(a);
130  float a2 = __high2float(a);
131  float b1 = __low2float(b);
132  float b2 = __high2float(b);
133  float r1 = a1 + b1;
134  float r2 = a2 + b2;
135  return __floats2half2_rn(r1, r2);
136 #endif
137 }
138 
139 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 psub<half2>(const half2& a, const half2& b) {
140 #if __CUDA_ARCH__ >= 530
141  return __hsub2(a, b);
142 #else
143  float a1 = __low2float(a);
144  float a2 = __high2float(a);
145  float b1 = __low2float(b);
146  float b2 = __high2float(b);
147  float r1 = a1 - b1;
148  float r2 = a2 - b2;
149  return __floats2half2_rn(r1, r2);
150 #endif
151 }
152 
153 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pnegate(const half2& a) {
154 #if __CUDA_ARCH__ >= 530
155  return __hneg2(a);
156 #else
157  float a1 = __low2float(a);
158  float a2 = __high2float(a);
159  return __floats2half2_rn(-a1, -a2);
160 #endif
161 }
162 
163 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pconj(const half2& a) { return a; }
164 
165 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pmul<half2>(const half2& a, const half2& b) {
166 #if __CUDA_ARCH__ >= 530
167  return __hmul2(a, b);
168 #else
169  float a1 = __low2float(a);
170  float a2 = __high2float(a);
171  float b1 = __low2float(b);
172  float b2 = __high2float(b);
173  float r1 = a1 * b1;
174  float r2 = a2 * b2;
175  return __floats2half2_rn(r1, r2);
176 #endif
177 }
178 
179 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pmadd<half2>(const half2& a, const half2& b, const half2& c) {
180 #if __CUDA_ARCH__ >= 530
181  return __hfma2(a, b, c);
182 #else
183  float a1 = __low2float(a);
184  float a2 = __high2float(a);
185  float b1 = __low2float(b);
186  float b2 = __high2float(b);
187  float c1 = __low2float(c);
188  float c2 = __high2float(c);
189  float r1 = a1 * b1 + c1;
190  float r2 = a2 * b2 + c2;
191  return __floats2half2_rn(r1, r2);
192 #endif
193 }
194 
195 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pdiv<half2>(const half2& a, const half2& b) {
196  float a1 = __low2float(a);
197  float a2 = __high2float(a);
198  float b1 = __low2float(b);
199  float b2 = __high2float(b);
200  float r1 = a1 / b1;
201  float r2 = a2 / b2;
202  return __floats2half2_rn(r1, r2);
203 }
204 
205 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pmin<half2>(const half2& a, const half2& b) {
206  float a1 = __low2float(a);
207  float a2 = __high2float(a);
208  float b1 = __low2float(b);
209  float b2 = __high2float(b);
210  __half r1 = a1 < b1 ? __low2half(a) : __low2half(b);
211  __half r2 = a2 < b2 ? __high2half(a) : __high2half(b);
212  return __halves2half2(r1, r2);
213 }
214 
215 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pmax<half2>(const half2& a, const half2& b) {
216  float a1 = __low2float(a);
217  float a2 = __high2float(a);
218  float b1 = __low2float(b);
219  float b2 = __high2float(b);
220  __half r1 = a1 > b1 ? __low2half(a) : __low2half(b);
221  __half r2 = a2 > b2 ? __high2half(a) : __high2half(b);
222  return __halves2half2(r1, r2);
223 }
224 
225 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::half predux<half2>(const half2& a) {
226 #if __CUDA_ARCH__ >= 530
227  return __hadd(__low2half(a), __high2half(a));
228 #else
229  float a1 = __low2float(a);
230  float a2 = __high2float(a);
231  return Eigen::half(half_impl::raw_uint16_to_half(__float2half_rn(a1 + a2)));
232 #endif
233 }
234 
235 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::half predux_max<half2>(const half2& a) {
236 #if __CUDA_ARCH__ >= 530
237  __half first = __low2half(a);
238  __half second = __high2half(a);
239  return __hgt(first, second) ? first : second;
240 #else
241  float a1 = __low2float(a);
242  float a2 = __high2float(a);
243  return a1 > a2 ? __low2half(a) : __high2half(a);
244 #endif
245 }
246 
247 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::half predux_min<half2>(const half2& a) {
248 #if __CUDA_ARCH__ >= 530
249  __half first = __low2half(a);
250  __half second = __high2half(a);
251  return __hlt(first, second) ? first : second;
252 #else
253  float a1 = __low2float(a);
254  float a2 = __high2float(a);
255  return a1 < a2 ? __low2half(a) : __high2half(a);
256 #endif
257 }
258 
259 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::half predux_mul<half2>(const half2& a) {
260 #if __CUDA_ARCH__ >= 530
261  return __hmul(__low2half(a), __high2half(a));
262 #else
263  float a1 = __low2float(a);
264  float a2 = __high2float(a);
265  return Eigen::half(half_impl::raw_uint16_to_half(__float2half_rn(a1 * a2)));
266 #endif
267 }
268 
269 #if defined __CUDACC_VER__ && __CUDACC_VER__ >= 80000 && defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 530
270 
271 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
272 half2 plog<half2>(const half2& a) {
273  return h2log(a);
274 }
275 
276 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
277 half2 pexp<half2>(const half2& a) {
278  return h2exp(a);
279 }
280 
281 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
282 half2 psqrt<half2>(const half2& a) {
283  return h2sqrt(a);
284 }
285 
286 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
287 half2 prsqrt<half2>(const half2& a) {
288  return h2rsqrt(a);
289 }
290 
291 #else
292 
293 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 plog<half2>(const half2& a) {
294  float a1 = __low2float(a);
295  float a2 = __high2float(a);
296  float r1 = logf(a1);
297  float r2 = logf(a2);
298  return __floats2half2_rn(r1, r2);
299 }
300 
301 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pexp<half2>(const half2& a) {
302  float a1 = __low2float(a);
303  float a2 = __high2float(a);
304  float r1 = expf(a1);
305  float r2 = expf(a2);
306  return __floats2half2_rn(r1, r2);
307 }
308 
309 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 psqrt<half2>(const half2& a) {
310  float a1 = __low2float(a);
311  float a2 = __high2float(a);
312  float r1 = sqrtf(a1);
313  float r2 = sqrtf(a2);
314  return __floats2half2_rn(r1, r2);
315 }
316 
317 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 prsqrt<half2>(const half2& a) {
318  float a1 = __low2float(a);
319  float a2 = __high2float(a);
320  float r1 = rsqrtf(a1);
321  float r2 = rsqrtf(a2);
322  return __floats2half2_rn(r1, r2);
323 }
324 
325 #endif
326 
327 #elif defined EIGEN_VECTORIZE_AVX
328 
329 typedef struct {
330  __m128i x;
331 } Packet8h;
332 
333 
334 template<> struct is_arithmetic<Packet8h> { enum { value = true }; };
335 
336 template <>
337 struct packet_traits<Eigen::half> : default_packet_traits {
338  typedef Packet8h type;
339  // There is no half-size packet for Packet8h.
340  typedef Packet8h half;
341  enum {
342  Vectorizable = 1,
343  AlignedOnScalar = 1,
344  size = 8,
345  HasHalfPacket = 0,
346  HasAdd = 0,
347  HasSub = 0,
348  HasMul = 0,
349  HasNegate = 0,
350  HasAbs = 0,
351  HasAbs2 = 0,
352  HasMin = 0,
353  HasMax = 0,
354  HasConj = 0,
355  HasSetLinear = 0,
356  HasDiv = 0,
357  HasSqrt = 0,
358  HasRsqrt = 0,
359  HasExp = 0,
360  HasLog = 0,
361  HasBlend = 0
362  };
363 };
364 
365 
366 template<> struct unpacket_traits<Packet8h> { typedef Eigen::half type; enum {size=8, alignment=Aligned16}; typedef Packet8h half; };
367 
368 template<> EIGEN_STRONG_INLINE Packet8h pset1<Packet8h>(const Eigen::half& from) {
369  Packet8h result;
370  result.x = _mm_set1_epi16(from.x);
371  return result;
372 }
373 
374 template<> EIGEN_STRONG_INLINE Eigen::half pfirst<Packet8h>(const Packet8h& from) {
375  return half_impl::raw_uint16_to_half(static_cast<unsigned short>(_mm_extract_epi16(from.x, 0)));
376 }
377 
378 template<> EIGEN_STRONG_INLINE Packet8h pload<Packet8h>(const Eigen::half* from) {
379  Packet8h result;
380  result.x = _mm_load_si128(reinterpret_cast<const __m128i*>(from));
381  return result;
382 }
383 
384 template<> EIGEN_STRONG_INLINE Packet8h ploadu<Packet8h>(const Eigen::half* from) {
385  Packet8h result;
386  result.x = _mm_loadu_si128(reinterpret_cast<const __m128i*>(from));
387  return result;
388 }
389 
390 template<> EIGEN_STRONG_INLINE void pstore<Eigen::half>(Eigen::half* to, const Packet8h& from) {
391  _mm_store_si128(reinterpret_cast<__m128i*>(to), from.x);
392 }
393 
394 template<> EIGEN_STRONG_INLINE void pstoreu<Eigen::half>(Eigen::half* to, const Packet8h& from) {
395  _mm_storeu_si128(reinterpret_cast<__m128i*>(to), from.x);
396 }
397 
398 template<> EIGEN_STRONG_INLINE Packet8h
399 ploadquad<Packet8h>(const Eigen::half* from) {
400  Packet8h result;
401  unsigned short a = from[0].x;
402  unsigned short b = from[1].x;
403  result.x = _mm_set_epi16(b, b, b, b, a, a, a, a);
404  return result;
405 }
406 
407 EIGEN_STRONG_INLINE Packet8f half2float(const Packet8h& a) {
408 #ifdef EIGEN_HAS_FP16_C
409  return _mm256_cvtph_ps(a.x);
410 #else
411  EIGEN_ALIGN32 Eigen::half aux[8];
412  pstore(aux, a);
413  float f0(aux[0]);
414  float f1(aux[1]);
415  float f2(aux[2]);
416  float f3(aux[3]);
417  float f4(aux[4]);
418  float f5(aux[5]);
419  float f6(aux[6]);
420  float f7(aux[7]);
421 
422  return _mm256_set_ps(f7, f6, f5, f4, f3, f2, f1, f0);
423 #endif
424 }
425 
426 EIGEN_STRONG_INLINE Packet8h float2half(const Packet8f& a) {
427 #ifdef EIGEN_HAS_FP16_C
428  Packet8h result;
429  result.x = _mm256_cvtps_ph(a, _MM_FROUND_TO_NEAREST_INT|_MM_FROUND_NO_EXC);
430  return result;
431 #else
432  EIGEN_ALIGN32 float aux[8];
433  pstore(aux, a);
434  Eigen::half h0(aux[0]);
435  Eigen::half h1(aux[1]);
436  Eigen::half h2(aux[2]);
437  Eigen::half h3(aux[3]);
438  Eigen::half h4(aux[4]);
439  Eigen::half h5(aux[5]);
440  Eigen::half h6(aux[6]);
441  Eigen::half h7(aux[7]);
442 
443  Packet8h result;
444  result.x = _mm_set_epi16(h7.x, h6.x, h5.x, h4.x, h3.x, h2.x, h1.x, h0.x);
445  return result;
446 #endif
447 }
448 
449 template<> EIGEN_STRONG_INLINE Packet8h pconj(const Packet8h& a) { return a; }
450 
451 template<> EIGEN_STRONG_INLINE Packet8h padd<Packet8h>(const Packet8h& a, const Packet8h& b) {
452  Packet8f af = half2float(a);
453  Packet8f bf = half2float(b);
454  Packet8f rf = padd(af, bf);
455  return float2half(rf);
456 }
457 
458 template<> EIGEN_STRONG_INLINE Packet8h pmul<Packet8h>(const Packet8h& a, const Packet8h& b) {
459  Packet8f af = half2float(a);
460  Packet8f bf = half2float(b);
461  Packet8f rf = pmul(af, bf);
462  return float2half(rf);
463 }
464 
465 template<> EIGEN_STRONG_INLINE Packet8h pgather<Eigen::half, Packet8h>(const Eigen::half* from, Index stride)
466 {
467  Packet8h result;
468  result.x = _mm_set_epi16(from[7*stride].x, from[6*stride].x, from[5*stride].x, from[4*stride].x, from[3*stride].x, from[2*stride].x, from[1*stride].x, from[0*stride].x);
469  return result;
470 }
471 
472 template<> EIGEN_STRONG_INLINE void pscatter<Eigen::half, Packet8h>(Eigen::half* to, const Packet8h& from, Index stride)
473 {
474  EIGEN_ALIGN32 Eigen::half aux[8];
475  pstore(aux, from);
476  to[stride*0].x = aux[0].x;
477  to[stride*1].x = aux[1].x;
478  to[stride*2].x = aux[2].x;
479  to[stride*3].x = aux[3].x;
480  to[stride*4].x = aux[4].x;
481  to[stride*5].x = aux[5].x;
482  to[stride*6].x = aux[6].x;
483  to[stride*7].x = aux[7].x;
484 }
485 
486 EIGEN_STRONG_INLINE void
487 ptranspose(PacketBlock<Packet8h,8>& kernel) {
488  __m128i a = kernel.packet[0].x;
489  __m128i b = kernel.packet[1].x;
490  __m128i c = kernel.packet[2].x;
491  __m128i d = kernel.packet[3].x;
492  __m128i e = kernel.packet[4].x;
493  __m128i f = kernel.packet[5].x;
494  __m128i g = kernel.packet[6].x;
495  __m128i h = kernel.packet[7].x;
496 
497  __m128i a03b03 = _mm_unpacklo_epi16(a, b);
498  __m128i c03d03 = _mm_unpacklo_epi16(c, d);
499  __m128i e03f03 = _mm_unpacklo_epi16(e, f);
500  __m128i g03h03 = _mm_unpacklo_epi16(g, h);
501  __m128i a47b47 = _mm_unpackhi_epi16(a, b);
502  __m128i c47d47 = _mm_unpackhi_epi16(c, d);
503  __m128i e47f47 = _mm_unpackhi_epi16(e, f);
504  __m128i g47h47 = _mm_unpackhi_epi16(g, h);
505 
506  __m128i a01b01c01d01 = _mm_unpacklo_epi32(a03b03, c03d03);
507  __m128i a23b23c23d23 = _mm_unpackhi_epi32(a03b03, c03d03);
508  __m128i e01f01g01h01 = _mm_unpacklo_epi32(e03f03, g03h03);
509  __m128i e23f23g23h23 = _mm_unpackhi_epi32(e03f03, g03h03);
510  __m128i a45b45c45d45 = _mm_unpacklo_epi32(a47b47, c47d47);
511  __m128i a67b67c67d67 = _mm_unpackhi_epi32(a47b47, c47d47);
512  __m128i e45f45g45h45 = _mm_unpacklo_epi32(e47f47, g47h47);
513  __m128i e67f67g67h67 = _mm_unpackhi_epi32(e47f47, g47h47);
514 
515  __m128i a0b0c0d0e0f0g0h0 = _mm_unpacklo_epi64(a01b01c01d01, e01f01g01h01);
516  __m128i a1b1c1d1e1f1g1h1 = _mm_unpackhi_epi64(a01b01c01d01, e01f01g01h01);
517  __m128i a2b2c2d2e2f2g2h2 = _mm_unpacklo_epi64(a23b23c23d23, e23f23g23h23);
518  __m128i a3b3c3d3e3f3g3h3 = _mm_unpackhi_epi64(a23b23c23d23, e23f23g23h23);
519  __m128i a4b4c4d4e4f4g4h4 = _mm_unpacklo_epi64(a45b45c45d45, e45f45g45h45);
520  __m128i a5b5c5d5e5f5g5h5 = _mm_unpackhi_epi64(a45b45c45d45, e45f45g45h45);
521  __m128i a6b6c6d6e6f6g6h6 = _mm_unpacklo_epi64(a67b67c67d67, e67f67g67h67);
522  __m128i a7b7c7d7e7f7g7h7 = _mm_unpackhi_epi64(a67b67c67d67, e67f67g67h67);
523 
524  kernel.packet[0].x = a0b0c0d0e0f0g0h0;
525  kernel.packet[1].x = a1b1c1d1e1f1g1h1;
526  kernel.packet[2].x = a2b2c2d2e2f2g2h2;
527  kernel.packet[3].x = a3b3c3d3e3f3g3h3;
528  kernel.packet[4].x = a4b4c4d4e4f4g4h4;
529  kernel.packet[5].x = a5b5c5d5e5f5g5h5;
530  kernel.packet[6].x = a6b6c6d6e6f6g6h6;
531  kernel.packet[7].x = a7b7c7d7e7f7g7h7;
532 }
533 
534 EIGEN_STRONG_INLINE void
535 ptranspose(PacketBlock<Packet8h,4>& kernel) {
536  EIGEN_ALIGN32 Eigen::half in[4][8];
537  pstore<Eigen::half>(in[0], kernel.packet[0]);
538  pstore<Eigen::half>(in[1], kernel.packet[1]);
539  pstore<Eigen::half>(in[2], kernel.packet[2]);
540  pstore<Eigen::half>(in[3], kernel.packet[3]);
541 
542  EIGEN_ALIGN32 Eigen::half out[4][8];
543 
544  for (int i = 0; i < 4; ++i) {
545  for (int j = 0; j < 4; ++j) {
546  out[i][j] = in[j][2*i];
547  }
548  for (int j = 0; j < 4; ++j) {
549  out[i][j+4] = in[j][2*i+1];
550  }
551  }
552 
553  kernel.packet[0] = pload<Packet8h>(out[0]);
554  kernel.packet[1] = pload<Packet8h>(out[1]);
555  kernel.packet[2] = pload<Packet8h>(out[2]);
556  kernel.packet[3] = pload<Packet8h>(out[3]);
557 }
558 
559 
560 // Disable the following code since it's broken on too many platforms / compilers.
561 //#elif defined(EIGEN_VECTORIZE_SSE) && (!EIGEN_ARCH_x86_64) && (!EIGEN_COMP_MSVC)
562 #elif 0
563 
564 typedef struct {
565  __m64 x;
566 } Packet4h;
567 
568 
569 template<> struct is_arithmetic<Packet4h> { enum { value = true }; };
570 
571 template <>
572 struct packet_traits<Eigen::half> : default_packet_traits {
573  typedef Packet4h type;
574  // There is no half-size packet for Packet4h.
575  typedef Packet4h half;
576  enum {
577  Vectorizable = 1,
578  AlignedOnScalar = 1,
579  size = 4,
580  HasHalfPacket = 0,
581  HasAdd = 0,
582  HasSub = 0,
583  HasMul = 0,
584  HasNegate = 0,
585  HasAbs = 0,
586  HasAbs2 = 0,
587  HasMin = 0,
588  HasMax = 0,
589  HasConj = 0,
590  HasSetLinear = 0,
591  HasDiv = 0,
592  HasSqrt = 0,
593  HasRsqrt = 0,
594  HasExp = 0,
595  HasLog = 0,
596  HasBlend = 0
597  };
598 };
599 
600 
601 template<> struct unpacket_traits<Packet4h> { typedef Eigen::half type; enum {size=4, alignment=Aligned16}; typedef Packet4h half; };
602 
603 template<> EIGEN_STRONG_INLINE Packet4h pset1<Packet4h>(const Eigen::half& from) {
604  Packet4h result;
605  result.x = _mm_set1_pi16(from.x);
606  return result;
607 }
608 
609 template<> EIGEN_STRONG_INLINE Eigen::half pfirst<Packet4h>(const Packet4h& from) {
610  return raw_uint16_to_half(static_cast<unsigned short>(_mm_cvtsi64_si32(from.x)));
611 }
612 
613 template<> EIGEN_STRONG_INLINE Packet4h pconj(const Packet4h& a) { return a; }
614 
615 template<> EIGEN_STRONG_INLINE Packet4h padd<Packet4h>(const Packet4h& a, const Packet4h& b) {
616  __int64_t a64 = _mm_cvtm64_si64(a.x);
617  __int64_t b64 = _mm_cvtm64_si64(b.x);
618 
619  Eigen::half h[4];
620 
621  Eigen::half ha = raw_uint16_to_half(static_cast<unsigned short>(a64));
622  Eigen::half hb = raw_uint16_to_half(static_cast<unsigned short>(b64));
623  h[0] = ha + hb;
624  ha = raw_uint16_to_half(static_cast<unsigned short>(a64 >> 16));
625  hb = raw_uint16_to_half(static_cast<unsigned short>(b64 >> 16));
626  h[1] = ha + hb;
627  ha = raw_uint16_to_half(static_cast<unsigned short>(a64 >> 32));
628  hb = raw_uint16_to_half(static_cast<unsigned short>(b64 >> 32));
629  h[2] = ha + hb;
630  ha = raw_uint16_to_half(static_cast<unsigned short>(a64 >> 48));
631  hb = raw_uint16_to_half(static_cast<unsigned short>(b64 >> 48));
632  h[3] = ha + hb;
633  Packet4h result;
634  result.x = _mm_set_pi16(h[3].x, h[2].x, h[1].x, h[0].x);
635  return result;
636 }
637 
638 template<> EIGEN_STRONG_INLINE Packet4h pmul<Packet4h>(const Packet4h& a, const Packet4h& b) {
639  __int64_t a64 = _mm_cvtm64_si64(a.x);
640  __int64_t b64 = _mm_cvtm64_si64(b.x);
641 
642  Eigen::half h[4];
643 
644  Eigen::half ha = raw_uint16_to_half(static_cast<unsigned short>(a64));
645  Eigen::half hb = raw_uint16_to_half(static_cast<unsigned short>(b64));
646  h[0] = ha * hb;
647  ha = raw_uint16_to_half(static_cast<unsigned short>(a64 >> 16));
648  hb = raw_uint16_to_half(static_cast<unsigned short>(b64 >> 16));
649  h[1] = ha * hb;
650  ha = raw_uint16_to_half(static_cast<unsigned short>(a64 >> 32));
651  hb = raw_uint16_to_half(static_cast<unsigned short>(b64 >> 32));
652  h[2] = ha * hb;
653  ha = raw_uint16_to_half(static_cast<unsigned short>(a64 >> 48));
654  hb = raw_uint16_to_half(static_cast<unsigned short>(b64 >> 48));
655  h[3] = ha * hb;
656  Packet4h result;
657  result.x = _mm_set_pi16(h[3].x, h[2].x, h[1].x, h[0].x);
658  return result;
659 }
660 
661 template<> EIGEN_STRONG_INLINE Packet4h pload<Packet4h>(const Eigen::half* from) {
662  Packet4h result;
663  result.x = _mm_cvtsi64_m64(*reinterpret_cast<const __int64_t*>(from));
664  return result;
665 }
666 
667 template<> EIGEN_STRONG_INLINE Packet4h ploadu<Packet4h>(const Eigen::half* from) {
668  Packet4h result;
669  result.x = _mm_cvtsi64_m64(*reinterpret_cast<const __int64_t*>(from));
670  return result;
671 }
672 
673 template<> EIGEN_STRONG_INLINE void pstore<Eigen::half>(Eigen::half* to, const Packet4h& from) {
674  __int64_t r = _mm_cvtm64_si64(from.x);
675  *(reinterpret_cast<__int64_t*>(to)) = r;
676 }
677 
678 template<> EIGEN_STRONG_INLINE void pstoreu<Eigen::half>(Eigen::half* to, const Packet4h& from) {
679  __int64_t r = _mm_cvtm64_si64(from.x);
680  *(reinterpret_cast<__int64_t*>(to)) = r;
681 }
682 
683 template<> EIGEN_STRONG_INLINE Packet4h
684 ploadquad<Packet4h>(const Eigen::half* from) {
685  return pset1<Packet4h>(*from);
686 }
687 
688 template<> EIGEN_STRONG_INLINE Packet4h pgather<Eigen::half, Packet4h>(const Eigen::half* from, Index stride)
689 {
690  Packet4h result;
691  result.x = _mm_set_pi16(from[3*stride].x, from[2*stride].x, from[1*stride].x, from[0*stride].x);
692  return result;
693 }
694 
695 template<> EIGEN_STRONG_INLINE void pscatter<Eigen::half, Packet4h>(Eigen::half* to, const Packet4h& from, Index stride)
696 {
697  __int64_t a = _mm_cvtm64_si64(from.x);
698  to[stride*0].x = static_cast<unsigned short>(a);
699  to[stride*1].x = static_cast<unsigned short>(a >> 16);
700  to[stride*2].x = static_cast<unsigned short>(a >> 32);
701  to[stride*3].x = static_cast<unsigned short>(a >> 48);
702 }
703 
704 EIGEN_STRONG_INLINE void
705 ptranspose(PacketBlock<Packet4h,4>& kernel) {
706  __m64 T0 = _mm_unpacklo_pi16(kernel.packet[0].x, kernel.packet[1].x);
707  __m64 T1 = _mm_unpacklo_pi16(kernel.packet[2].x, kernel.packet[3].x);
708  __m64 T2 = _mm_unpackhi_pi16(kernel.packet[0].x, kernel.packet[1].x);
709  __m64 T3 = _mm_unpackhi_pi16(kernel.packet[2].x, kernel.packet[3].x);
710 
711  kernel.packet[0].x = _mm_unpacklo_pi32(T0, T1);
712  kernel.packet[1].x = _mm_unpackhi_pi32(T0, T1);
713  kernel.packet[2].x = _mm_unpacklo_pi32(T2, T3);
714  kernel.packet[3].x = _mm_unpackhi_pi32(T2, T3);
715 }
716 
717 #endif
718 
719 }
720 }
721 
722 #endif // EIGEN_PACKET_MATH_HALF_CUDA_H
Definition: Constants.h:230
Namespace containing all symbols from the Eigen library.
Definition: Core:271
EIGEN_DEFAULT_DENSE_INDEX_TYPE Index
The Index type as used for the API.
Definition: XprHelper.h:35
Definition: Eigen_Colamd.h:50