12b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang// This file is part of Eigen, a lightweight C++ template library 22b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang// for linear algebra. 32b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang// 42b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang// Copyright (C) 2014 Benoit Steiner <benoit.steiner.goog@gmail.com> 52b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang// 62b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang// This Source Code Form is subject to the terms of the Mozilla 72b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang// Public License v. 2.0. If a copy of the MPL was not distributed 82b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang// with this file, You can obtain one at http://mozilla.org/MPL/2.0/. 92b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 102b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang#ifndef EIGEN_PACKET_MATH_CUDA_H 112b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang#define EIGEN_PACKET_MATH_CUDA_H 122b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 132b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangnamespace Eigen { 142b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 152b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangnamespace internal { 162b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 172b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang// Make sure this is only available when targeting a GPU: we don't want to 182b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang// introduce conflicts between these packet_traits definitions and the ones 192b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang// we'll use on the host side (SSE, AVX, ...) 202b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang#if defined(__CUDACC__) && defined(EIGEN_USE_GPU) 212b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> struct is_arithmetic<float4> { enum { value = true }; }; 222b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> struct is_arithmetic<double2> { enum { value = true }; }; 232b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 242b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> struct packet_traits<float> : default_packet_traits 252b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang{ 262b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef float4 type; 272b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef float4 half; 282b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang enum { 292b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang Vectorizable = 1, 302b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang AlignedOnScalar = 1, 312b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang size=4, 322b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang HasHalfPacket = 0, 332b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 342b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang HasDiv = 1, 352b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang HasSin = 0, 362b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang HasCos = 0, 372b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang HasLog = 1, 382b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang HasExp = 1, 392b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang HasSqrt = 1, 402b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang HasRsqrt = 1, 412b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang HasLGamma = 1, 422b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang HasDiGamma = 1, 432b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang HasZeta = 1, 442b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang HasPolygamma = 1, 452b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang HasErf = 1, 462b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang HasErfc = 1, 472b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang HasIGamma = 1, 482b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang HasIGammac = 1, 492b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang HasBetaInc = 1, 502b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 512b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang HasBlend = 0, 522b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang }; 532b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang}; 542b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 552b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> struct packet_traits<double> : default_packet_traits 562b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang{ 572b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef double2 type; 582b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang typedef double2 half; 592b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang enum { 602b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang Vectorizable = 1, 612b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang AlignedOnScalar = 1, 622b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang size=2, 632b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang HasHalfPacket = 0, 642b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 652b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang HasDiv = 1, 662b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang HasLog = 1, 672b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang HasExp = 1, 682b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang HasSqrt = 1, 692b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang HasRsqrt = 1, 702b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang HasLGamma = 1, 712b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang HasDiGamma = 1, 722b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang HasZeta = 1, 732b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang HasPolygamma = 1, 742b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang HasErf = 1, 752b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang HasErfc = 1, 762b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang HasIGamma = 1, 772b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang HasIGammac = 1, 782b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang HasBetaInc = 1, 792b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 802b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang HasBlend = 0, 812b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang }; 822b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang}; 832b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 842b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 852b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> struct unpacket_traits<float4> { typedef float type; enum {size=4, alignment=Aligned16}; typedef float4 half; }; 862b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> struct unpacket_traits<double2> { typedef double type; enum {size=2, alignment=Aligned16}; typedef double2 half; }; 872b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 882b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pset1<float4>(const float& from) { 892b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return make_float4(from, from, from, from); 902b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 912b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 pset1<double2>(const double& from) { 922b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return make_double2(from, from); 932b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 942b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 952b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 962b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 plset<float4>(const float& a) { 972b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return make_float4(a, a+1, a+2, a+3); 982b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 992b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 plset<double2>(const double& a) { 1002b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return make_double2(a, a+1); 1012b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 1022b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 1032b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 padd<float4>(const float4& a, const float4& b) { 1042b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return make_float4(a.x+b.x, a.y+b.y, a.z+b.z, a.w+b.w); 1052b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 1062b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 padd<double2>(const double2& a, const double2& b) { 1072b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return make_double2(a.x+b.x, a.y+b.y); 1082b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 1092b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 1102b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 psub<float4>(const float4& a, const float4& b) { 1112b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return make_float4(a.x-b.x, a.y-b.y, a.z-b.z, a.w-b.w); 1122b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 1132b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 psub<double2>(const double2& a, const double2& b) { 1142b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return make_double2(a.x-b.x, a.y-b.y); 1152b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 1162b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 1172b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pnegate(const float4& a) { 1182b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return make_float4(-a.x, -a.y, -a.z, -a.w); 1192b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 1202b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 pnegate(const double2& a) { 1212b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return make_double2(-a.x, -a.y); 1222b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 1232b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 1242b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pconj(const float4& a) { return a; } 1252b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 pconj(const double2& a) { return a; } 1262b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 1272b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pmul<float4>(const float4& a, const float4& b) { 1282b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return make_float4(a.x*b.x, a.y*b.y, a.z*b.z, a.w*b.w); 1292b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 1302b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 pmul<double2>(const double2& a, const double2& b) { 1312b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return make_double2(a.x*b.x, a.y*b.y); 1322b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 1332b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 1342b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pdiv<float4>(const float4& a, const float4& b) { 1352b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return make_float4(a.x/b.x, a.y/b.y, a.z/b.z, a.w/b.w); 1362b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 1372b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 pdiv<double2>(const double2& a, const double2& b) { 1382b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return make_double2(a.x/b.x, a.y/b.y); 1392b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 1402b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 1412b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pmin<float4>(const float4& a, const float4& b) { 1422b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return make_float4(fminf(a.x, b.x), fminf(a.y, b.y), fminf(a.z, b.z), fminf(a.w, b.w)); 1432b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 1442b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 pmin<double2>(const double2& a, const double2& b) { 1452b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return make_double2(fmin(a.x, b.x), fmin(a.y, b.y)); 1462b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 1472b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 1482b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pmax<float4>(const float4& a, const float4& b) { 1492b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return make_float4(fmaxf(a.x, b.x), fmaxf(a.y, b.y), fmaxf(a.z, b.z), fmaxf(a.w, b.w)); 1502b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 1512b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 pmax<double2>(const double2& a, const double2& b) { 1522b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return make_double2(fmax(a.x, b.x), fmax(a.y, b.y)); 1532b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 1542b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 1552b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pload<float4>(const float* from) { 1562b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return *reinterpret_cast<const float4*>(from); 1572b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 1582b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 1592b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 pload<double2>(const double* from) { 1602b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return *reinterpret_cast<const double2*>(from); 1612b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 1622b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 1632b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 ploadu<float4>(const float* from) { 1642b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return make_float4(from[0], from[1], from[2], from[3]); 1652b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 1662b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 ploadu<double2>(const double* from) { 1672b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return make_double2(from[0], from[1]); 1682b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 1692b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 1702b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_STRONG_INLINE float4 ploaddup<float4>(const float* from) { 1712b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return make_float4(from[0], from[0], from[1], from[1]); 1722b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 1732b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_STRONG_INLINE double2 ploaddup<double2>(const double* from) { 1742b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return make_double2(from[0], from[0]); 1752b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 1762b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 1772b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void pstore<float>(float* to, const float4& from) { 1782b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang *reinterpret_cast<float4*>(to) = from; 1792b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 1802b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 1812b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void pstore<double>(double* to, const double2& from) { 1822b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang *reinterpret_cast<double2*>(to) = from; 1832b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 1842b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 1852b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void pstoreu<float>(float* to, const float4& from) { 1862b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang to[0] = from.x; 1872b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang to[1] = from.y; 1882b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang to[2] = from.z; 1892b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang to[3] = from.w; 1902b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 1912b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 1922b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void pstoreu<double>(double* to, const double2& from) { 1932b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang to[0] = from.x; 1942b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang to[1] = from.y; 1952b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 1962b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 1972b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> 1982b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao WangEIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float4 ploadt_ro<float4, Aligned>(const float* from) { 1992b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 350 2002b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return __ldg((const float4*)from); 2012b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang#else 2022b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return make_float4(from[0], from[1], from[2], from[3]); 2032b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang#endif 2042b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 2052b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> 2062b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao WangEIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double2 ploadt_ro<double2, Aligned>(const double* from) { 2072b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 350 2082b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return __ldg((const double2*)from); 2092b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang#else 2102b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return make_double2(from[0], from[1]); 2112b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang#endif 2122b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 2132b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2142b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> 2152b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao WangEIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float4 ploadt_ro<float4, Unaligned>(const float* from) { 2162b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 350 2172b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return make_float4(__ldg(from+0), __ldg(from+1), __ldg(from+2), __ldg(from+3)); 2182b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang#else 2192b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return make_float4(from[0], from[1], from[2], from[3]); 2202b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang#endif 2212b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 2222b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> 2232b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao WangEIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double2 ploadt_ro<double2, Unaligned>(const double* from) { 2242b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 350 2252b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return make_double2(__ldg(from+0), __ldg(from+1)); 2262b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang#else 2272b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return make_double2(from[0], from[1]); 2282b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang#endif 2292b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 2302b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2312b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC inline float4 pgather<float, float4>(const float* from, Index stride) { 2322b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return make_float4(from[0*stride], from[1*stride], from[2*stride], from[3*stride]); 2332b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 2342b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2352b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC inline double2 pgather<double, double2>(const double* from, Index stride) { 2362b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return make_double2(from[0*stride], from[1*stride]); 2372b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 2382b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2392b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC inline void pscatter<float, float4>(float* to, const float4& from, Index stride) { 2402b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang to[stride*0] = from.x; 2412b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang to[stride*1] = from.y; 2422b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang to[stride*2] = from.z; 2432b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang to[stride*3] = from.w; 2442b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 2452b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC inline void pscatter<double, double2>(double* to, const double2& from, Index stride) { 2462b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang to[stride*0] = from.x; 2472b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang to[stride*1] = from.y; 2482b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 2492b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2502b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC inline float pfirst<float4>(const float4& a) { 2512b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return a.x; 2522b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 2532b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC inline double pfirst<double2>(const double2& a) { 2542b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return a.x; 2552b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 2562b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2572b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC inline float predux<float4>(const float4& a) { 2582b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return a.x + a.y + a.z + a.w; 2592b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 2602b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC inline double predux<double2>(const double2& a) { 2612b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return a.x + a.y; 2622b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 2632b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2642b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC inline float predux_max<float4>(const float4& a) { 2652b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return fmaxf(fmaxf(a.x, a.y), fmaxf(a.z, a.w)); 2662b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 2672b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC inline double predux_max<double2>(const double2& a) { 2682b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return fmax(a.x, a.y); 2692b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 2702b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2712b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC inline float predux_min<float4>(const float4& a) { 2722b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return fminf(fminf(a.x, a.y), fminf(a.z, a.w)); 2732b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 2742b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC inline double predux_min<double2>(const double2& a) { 2752b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return fmin(a.x, a.y); 2762b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 2772b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2782b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC inline float predux_mul<float4>(const float4& a) { 2792b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return a.x * a.y * a.z * a.w; 2802b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 2812b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC inline double predux_mul<double2>(const double2& a) { 2822b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return a.x * a.y; 2832b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 2842b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2852b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC inline float4 pabs<float4>(const float4& a) { 2862b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return make_float4(fabsf(a.x), fabsf(a.y), fabsf(a.z), fabsf(a.w)); 2872b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 2882b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangtemplate<> EIGEN_DEVICE_FUNC inline double2 pabs<double2>(const double2& a) { 2892b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang return make_double2(fabs(a.x), fabs(a.y)); 2902b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 2912b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2922b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao WangEIGEN_DEVICE_FUNC inline void 2932b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangptranspose(PacketBlock<float4,4>& kernel) { 294eda03298de395cf6217486971e6529f92da8da79Miao Wang float tmp = kernel.packet[0].y; 2952b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang kernel.packet[0].y = kernel.packet[1].x; 2962b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang kernel.packet[1].x = tmp; 2972b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 2982b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang tmp = kernel.packet[0].z; 2992b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang kernel.packet[0].z = kernel.packet[2].x; 3002b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang kernel.packet[2].x = tmp; 3012b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 3022b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang tmp = kernel.packet[0].w; 3032b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang kernel.packet[0].w = kernel.packet[3].x; 3042b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang kernel.packet[3].x = tmp; 3052b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 3062b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang tmp = kernel.packet[1].z; 3072b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang kernel.packet[1].z = kernel.packet[2].y; 3082b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang kernel.packet[2].y = tmp; 3092b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 3102b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang tmp = kernel.packet[1].w; 3112b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang kernel.packet[1].w = kernel.packet[3].y; 3122b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang kernel.packet[3].y = tmp; 3132b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 3142b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang tmp = kernel.packet[2].w; 3152b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang kernel.packet[2].w = kernel.packet[3].z; 3162b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang kernel.packet[3].z = tmp; 3172b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 3182b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 3192b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao WangEIGEN_DEVICE_FUNC inline void 3202b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wangptranspose(PacketBlock<double2,2>& kernel) { 3212b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang double tmp = kernel.packet[0].y; 3222b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang kernel.packet[0].y = kernel.packet[1].x; 3232b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang kernel.packet[1].x = tmp; 3242b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} 3252b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 3262b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang#endif 3272b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 3282b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} // end namespace internal 3292b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 3302b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang} // end namespace Eigen 3312b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 3322b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang 3332b8756b6f1de65d3f8bffab45be6c44ceb7411fcMiao Wang#endif // EIGEN_PACKET_MATH_CUDA_H 334