10 #ifndef EIGEN_PACKET_MATH_HALF_CUDA_H 11 #define EIGEN_PACKET_MATH_HALF_CUDA_H 18 #if defined(EIGEN_HAS_CUDA_FP16) && defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 300 20 template<>
struct is_arithmetic<half2> {
enum { value =
true }; };
22 template<>
struct packet_traits<
Eigen::half> : default_packet_traits
41 template<>
struct unpacket_traits<half2> {
typedef Eigen::half type;
enum {size=2, alignment=
Aligned16};
typedef half2 half; };
43 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pset1<half2>(
const Eigen::half& from) {
44 return __half2half2(from);
47 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pload<half2>(
const Eigen::half* from) {
48 return *
reinterpret_cast<const half2*
>(from);
51 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 ploadu<half2>(
const Eigen::half* from) {
52 return __halves2half2(from[0], from[1]);
55 template<> EIGEN_STRONG_INLINE half2 ploaddup<half2>(
const Eigen::half* from) {
56 return __halves2half2(from[0], from[0]);
59 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
void pstore<Eigen::half>(Eigen::half* to,
const half2& from) {
60 *
reinterpret_cast<half2*
>(to) = from;
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);
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);
73 return __halves2half2(*(from+0), *(from+1));
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));
82 return __halves2half2(*(from+0), *(from+1));
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]);
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);
95 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::half pfirst<half2>(
const half2& a) {
99 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pabs<half2>(
const half2& a) {
101 result.x = a.x & 0x7FFF7FFF;
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);
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)));
120 float f = __half2float(a) + 1.0f;
121 return __halves2half2(a, __float2half(f));
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);
129 float a1 = __low2float(a);
130 float a2 = __high2float(a);
131 float b1 = __low2float(b);
132 float b2 = __high2float(b);
135 return __floats2half2_rn(r1, r2);
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);
143 float a1 = __low2float(a);
144 float a2 = __high2float(a);
145 float b1 = __low2float(b);
146 float b2 = __high2float(b);
149 return __floats2half2_rn(r1, r2);
153 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pnegate(
const half2& a) {
154 #if __CUDA_ARCH__ >= 530 157 float a1 = __low2float(a);
158 float a2 = __high2float(a);
159 return __floats2half2_rn(-a1, -a2);
163 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pconj(
const half2& a) {
return a; }
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);
169 float a1 = __low2float(a);
170 float a2 = __high2float(a);
171 float b1 = __low2float(b);
172 float b2 = __high2float(b);
175 return __floats2half2_rn(r1, r2);
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);
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);
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);
202 return __floats2half2_rn(r1, r2);
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);
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);
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));
229 float a1 = __low2float(a);
230 float a2 = __high2float(a);
231 return Eigen::half(half_impl::raw_uint16_to_half(__float2half_rn(a1 + a2)));
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;
241 float a1 = __low2float(a);
242 float a2 = __high2float(a);
243 return a1 > a2 ? __low2half(a) : __high2half(a);
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;
253 float a1 = __low2float(a);
254 float a2 = __high2float(a);
255 return a1 < a2 ? __low2half(a) : __high2half(a);
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));
263 float a1 = __low2float(a);
264 float a2 = __high2float(a);
265 return Eigen::half(half_impl::raw_uint16_to_half(__float2half_rn(a1 * a2)));
269 #if defined __CUDACC_VER__ && __CUDACC_VER__ >= 80000 && defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 530 271 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
272 half2 plog<half2>(
const half2& a) {
276 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
277 half2 pexp<half2>(
const half2& a) {
281 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
282 half2 psqrt<half2>(
const half2& a) {
286 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
287 half2 prsqrt<half2>(
const half2& a) {
293 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 plog<half2>(
const half2& a) {
294 float a1 = __low2float(a);
295 float a2 = __high2float(a);
298 return __floats2half2_rn(r1, r2);
301 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pexp<half2>(
const half2& a) {
302 float a1 = __low2float(a);
303 float a2 = __high2float(a);
306 return __floats2half2_rn(r1, r2);
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);
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);
327 #elif defined EIGEN_VECTORIZE_AVX 334 template<>
struct is_arithmetic<Packet8h> {
enum { value =
true }; };
337 struct packet_traits<
Eigen::half> : default_packet_traits {
338 typedef Packet8h type;
340 typedef Packet8h half;
366 template<>
struct unpacket_traits<Packet8h> {
typedef Eigen::half type;
enum {size=8, alignment=
Aligned16};
typedef Packet8h half; };
368 template<> EIGEN_STRONG_INLINE Packet8h pset1<Packet8h>(
const Eigen::half& from) {
370 result.x = _mm_set1_epi16(from.x);
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)));
378 template<> EIGEN_STRONG_INLINE Packet8h pload<Packet8h>(
const Eigen::half* from) {
380 result.x = _mm_load_si128(reinterpret_cast<const __m128i*>(from));
384 template<> EIGEN_STRONG_INLINE Packet8h ploadu<Packet8h>(
const Eigen::half* from) {
386 result.x = _mm_loadu_si128(reinterpret_cast<const __m128i*>(from));
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);
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);
398 template<> EIGEN_STRONG_INLINE Packet8h
399 ploadquad<Packet8h>(
const Eigen::half* from) {
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);
407 EIGEN_STRONG_INLINE Packet8f half2float(
const Packet8h& a) {
408 #ifdef EIGEN_HAS_FP16_C 409 return _mm256_cvtph_ps(a.x);
411 EIGEN_ALIGN32 Eigen::half aux[8];
422 return _mm256_set_ps(f7, f6, f5, f4, f3, f2, f1, f0);
426 EIGEN_STRONG_INLINE Packet8h float2half(
const Packet8f& a) {
427 #ifdef EIGEN_HAS_FP16_C 429 result.x = _mm256_cvtps_ph(a, _MM_FROUND_TO_NEAREST_INT|_MM_FROUND_NO_EXC);
432 EIGEN_ALIGN32
float aux[8];
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]);
444 result.x = _mm_set_epi16(h7.x, h6.x, h5.x, h4.x, h3.x, h2.x, h1.x, h0.x);
449 template<> EIGEN_STRONG_INLINE Packet8h pconj(
const Packet8h& a) {
return a; }
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);
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);
465 template<> EIGEN_STRONG_INLINE Packet8h pgather<Eigen::half, Packet8h>(
const Eigen::half* from,
Index stride)
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);
472 template<> EIGEN_STRONG_INLINE
void pscatter<Eigen::half, Packet8h>(Eigen::half* to,
const Packet8h& from,
Index stride)
474 EIGEN_ALIGN32 Eigen::half aux[8];
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;
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;
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);
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);
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);
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;
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]);
542 EIGEN_ALIGN32 Eigen::half out[4][8];
544 for (
int i = 0; i < 4; ++i) {
545 for (
int j = 0; j < 4; ++j) {
546 out[i][j] = in[j][2*i];
548 for (
int j = 0; j < 4; ++j) {
549 out[i][j+4] = in[j][2*i+1];
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]);
569 template<>
struct is_arithmetic<Packet4h> {
enum { value =
true }; };
572 struct packet_traits<
Eigen::half> : default_packet_traits {
573 typedef Packet4h type;
575 typedef Packet4h half;
601 template<>
struct unpacket_traits<Packet4h> {
typedef Eigen::half type;
enum {size=4, alignment=
Aligned16};
typedef Packet4h half; };
603 template<> EIGEN_STRONG_INLINE Packet4h pset1<Packet4h>(
const Eigen::half& from) {
605 result.x = _mm_set1_pi16(from.x);
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)));
613 template<> EIGEN_STRONG_INLINE Packet4h pconj(
const Packet4h& a) {
return a; }
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);
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));
624 ha = raw_uint16_to_half(static_cast<unsigned short>(a64 >> 16));
625 hb = raw_uint16_to_half(static_cast<unsigned short>(b64 >> 16));
627 ha = raw_uint16_to_half(static_cast<unsigned short>(a64 >> 32));
628 hb = raw_uint16_to_half(static_cast<unsigned short>(b64 >> 32));
630 ha = raw_uint16_to_half(static_cast<unsigned short>(a64 >> 48));
631 hb = raw_uint16_to_half(static_cast<unsigned short>(b64 >> 48));
634 result.x = _mm_set_pi16(h[3].x, h[2].x, h[1].x, h[0].x);
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);
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));
647 ha = raw_uint16_to_half(static_cast<unsigned short>(a64 >> 16));
648 hb = raw_uint16_to_half(static_cast<unsigned short>(b64 >> 16));
650 ha = raw_uint16_to_half(static_cast<unsigned short>(a64 >> 32));
651 hb = raw_uint16_to_half(static_cast<unsigned short>(b64 >> 32));
653 ha = raw_uint16_to_half(static_cast<unsigned short>(a64 >> 48));
654 hb = raw_uint16_to_half(static_cast<unsigned short>(b64 >> 48));
657 result.x = _mm_set_pi16(h[3].x, h[2].x, h[1].x, h[0].x);
661 template<> EIGEN_STRONG_INLINE Packet4h pload<Packet4h>(
const Eigen::half* from) {
663 result.x = _mm_cvtsi64_m64(*reinterpret_cast<const __int64_t*>(from));
667 template<> EIGEN_STRONG_INLINE Packet4h ploadu<Packet4h>(
const Eigen::half* from) {
669 result.x = _mm_cvtsi64_m64(*reinterpret_cast<const __int64_t*>(from));
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;
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;
683 template<> EIGEN_STRONG_INLINE Packet4h
684 ploadquad<Packet4h>(
const Eigen::half* from) {
685 return pset1<Packet4h>(*from);
688 template<> EIGEN_STRONG_INLINE Packet4h pgather<Eigen::half, Packet4h>(
const Eigen::half* from,
Index stride)
691 result.x = _mm_set_pi16(from[3*stride].x, from[2*stride].x, from[1*stride].x, from[0*stride].x);
695 template<> EIGEN_STRONG_INLINE
void pscatter<Eigen::half, Packet4h>(Eigen::half* to,
const Packet4h& from,
Index stride)
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);
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);
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);
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