Eigen  3.2.91
CUDA/PacketMath.h
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_CUDA_H
11 #define EIGEN_PACKET_MATH_CUDA_H
12 
13 namespace Eigen {
14 
15 namespace internal {
16 
17 // Make sure this is only available when targeting a GPU: we don't want to
18 // introduce conflicts between these packet_traits definitions and the ones
19 // we'll use on the host side (SSE, AVX, ...)
20 #if defined(__CUDACC__) && defined(EIGEN_USE_GPU)
21 template<> struct is_arithmetic<float4> { enum { value = true }; };
22 template<> struct is_arithmetic<double2> { enum { value = true }; };
23 
24 
25 template<> struct packet_traits<float> : default_packet_traits
26 {
27  typedef float4 type;
28  typedef float4 half;
29  enum {
30  Vectorizable = 1,
31  AlignedOnScalar = 1,
32  size=4,
33  HasHalfPacket = 0,
34 
35  HasDiv = 1,
36  HasSin = 0,
37  HasCos = 0,
38  HasLog = 1,
39  HasExp = 1,
40  HasSqrt = 1,
41  HasRsqrt = 1,
42 
43  HasBlend = 0,
44  };
45 };
46 
47 template<> struct packet_traits<double> : default_packet_traits
48 {
49  typedef double2 type;
50  typedef double2 half;
51  enum {
52  Vectorizable = 1,
53  AlignedOnScalar = 1,
54  size=2,
55  HasHalfPacket = 0,
56 
57  HasDiv = 1,
58  HasLog = 1,
59  HasExp = 1,
60  HasSqrt = 1,
61  HasRsqrt = 1,
62 
63  HasBlend = 0,
64  };
65 };
66 
67 
68 template<> struct unpacket_traits<float4> { typedef float type; enum {size=4, alignment=Aligned16}; typedef float4 half; };
69 template<> struct unpacket_traits<double2> { typedef double type; enum {size=2, alignment=Aligned16}; typedef double2 half; };
70 
71 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pset1<float4>(const float& from) {
72  return make_float4(from, from, from, from);
73 }
74 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 pset1<double2>(const double& from) {
75  return make_double2(from, from);
76 }
77 
78 
79 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 plset<float4>(const float& a) {
80  return make_float4(a, a+1, a+2, a+3);
81 }
82 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 plset<double2>(const double& a) {
83  return make_double2(a, a+1);
84 }
85 
86 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 padd<float4>(const float4& a, const float4& b) {
87  return make_float4(a.x+b.x, a.y+b.y, a.z+b.z, a.w+b.w);
88 }
89 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 padd<double2>(const double2& a, const double2& b) {
90  return make_double2(a.x+b.x, a.y+b.y);
91 }
92 
93 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 psub<float4>(const float4& a, const float4& b) {
94  return make_float4(a.x-b.x, a.y-b.y, a.z-b.z, a.w-b.w);
95 }
96 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 psub<double2>(const double2& a, const double2& b) {
97  return make_double2(a.x-b.x, a.y-b.y);
98 }
99 
100 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pnegate(const float4& a) {
101  return make_float4(-a.x, -a.y, -a.z, -a.w);
102 }
103 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 pnegate(const double2& a) {
104  return make_double2(-a.x, -a.y);
105 }
106 
107 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pconj(const float4& a) { return a; }
108 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 pconj(const double2& a) { return a; }
109 
110 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pmul<float4>(const float4& a, const float4& b) {
111  return make_float4(a.x*b.x, a.y*b.y, a.z*b.z, a.w*b.w);
112 }
113 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 pmul<double2>(const double2& a, const double2& b) {
114  return make_double2(a.x*b.x, a.y*b.y);
115 }
116 
117 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pdiv<float4>(const float4& a, const float4& b) {
118  return make_float4(a.x/b.x, a.y/b.y, a.z/b.z, a.w/b.w);
119 }
120 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 pdiv<double2>(const double2& a, const double2& b) {
121  return make_double2(a.x/b.x, a.y/b.y);
122 }
123 
124 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pmin<float4>(const float4& a, const float4& b) {
125  return make_float4(fminf(a.x, b.x), fminf(a.y, b.y), fminf(a.z, b.z), fminf(a.w, b.w));
126 }
127 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 pmin<double2>(const double2& a, const double2& b) {
128  return make_double2(fmin(a.x, b.x), fmin(a.y, b.y));
129 }
130 
131 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pmax<float4>(const float4& a, const float4& b) {
132  return make_float4(fmaxf(a.x, b.x), fmaxf(a.y, b.y), fmaxf(a.z, b.z), fmaxf(a.w, b.w));
133 }
134 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 pmax<double2>(const double2& a, const double2& b) {
135  return make_double2(fmax(a.x, b.x), fmax(a.y, b.y));
136 }
137 
138 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pload<float4>(const float* from) {
139  return *reinterpret_cast<const float4*>(from);
140 }
141 
142 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 pload<double2>(const double* from) {
143  return *reinterpret_cast<const double2*>(from);
144 }
145 
146 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 ploadu<float4>(const float* from) {
147  return make_float4(from[0], from[1], from[2], from[3]);
148 }
149 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 ploadu<double2>(const double* from) {
150  return make_double2(from[0], from[1]);
151 }
152 
153 template<> EIGEN_STRONG_INLINE float4 ploaddup<float4>(const float* from) {
154  return make_float4(from[0], from[0], from[1], from[1]);
155 }
156 template<> EIGEN_STRONG_INLINE double2 ploaddup<double2>(const double* from) {
157  return make_double2(from[0], from[0]);
158 }
159 
160 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void pstore<float>(float* to, const float4& from) {
161  *reinterpret_cast<float4*>(to) = from;
162 }
163 
164 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void pstore<double>(double* to, const double2& from) {
165  *reinterpret_cast<double2*>(to) = from;
166 }
167 
168 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void pstoreu<float>(float* to, const float4& from) {
169  to[0] = from.x;
170  to[1] = from.y;
171  to[2] = from.z;
172  to[3] = from.w;
173 }
174 
175 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void pstoreu<double>(double* to, const double2& from) {
176  to[0] = from.x;
177  to[1] = from.y;
178 }
179 
180 #ifdef __CUDA_ARCH__
181 template<>
182 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float4 ploadt_ro<float4, Aligned>(const float* from) {
183  return __ldg((const float4*)from);
184 }
185 template<>
186 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double2 ploadt_ro<double2, Aligned>(const double* from) {
187  return __ldg((const double2*)from);
188 }
189 
190 template<>
191 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float4 ploadt_ro<float4, Unaligned>(const float* from) {
192  return make_float4(__ldg(from+0), __ldg(from+1), __ldg(from+2), __ldg(from+3));
193 }
194 template<>
195 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double2 ploadt_ro<double2, Unaligned>(const double* from) {
196  return make_double2(__ldg(from+0), __ldg(from+1));
197 }
198 #endif
199 
200 template<> EIGEN_DEVICE_FUNC inline float4 pgather<float, float4>(const float* from, Index stride) {
201  return make_float4(from[0*stride], from[1*stride], from[2*stride], from[3*stride]);
202 }
203 
204 template<> EIGEN_DEVICE_FUNC inline double2 pgather<double, double2>(const double* from, Index stride) {
205  return make_double2(from[0*stride], from[1*stride]);
206 }
207 
208 template<> EIGEN_DEVICE_FUNC inline void pscatter<float, float4>(float* to, const float4& from, Index stride) {
209  to[stride*0] = from.x;
210  to[stride*1] = from.y;
211  to[stride*2] = from.z;
212  to[stride*3] = from.w;
213 }
214 template<> EIGEN_DEVICE_FUNC inline void pscatter<double, double2>(double* to, const double2& from, Index stride) {
215  to[stride*0] = from.x;
216  to[stride*1] = from.y;
217 }
218 
219 template<> EIGEN_DEVICE_FUNC inline float pfirst<float4>(const float4& a) {
220  return a.x;
221 }
222 template<> EIGEN_DEVICE_FUNC inline double pfirst<double2>(const double2& a) {
223  return a.x;
224 }
225 
226 template<> EIGEN_DEVICE_FUNC inline float predux<float4>(const float4& a) {
227  return a.x + a.y + a.z + a.w;
228 }
229 template<> EIGEN_DEVICE_FUNC inline double predux<double2>(const double2& a) {
230  return a.x + a.y;
231 }
232 
233 template<> EIGEN_DEVICE_FUNC inline float predux_max<float4>(const float4& a) {
234  return fmaxf(fmaxf(a.x, a.y), fmaxf(a.z, a.w));
235 }
236 template<> EIGEN_DEVICE_FUNC inline double predux_max<double2>(const double2& a) {
237  return fmax(a.x, a.y);
238 }
239 
240 template<> EIGEN_DEVICE_FUNC inline float predux_min<float4>(const float4& a) {
241  return fminf(fminf(a.x, a.y), fminf(a.z, a.w));
242 }
243 template<> EIGEN_DEVICE_FUNC inline double predux_min<double2>(const double2& a) {
244  return fmin(a.x, a.y);
245 }
246 
247 template<> EIGEN_DEVICE_FUNC inline float4 pabs<float4>(const float4& a) {
248  return make_float4(fabsf(a.x), fabsf(a.y), fabsf(a.z), fabsf(a.w));
249 }
250 template<> EIGEN_DEVICE_FUNC inline double2 pabs<double2>(const double2& a) {
251  return make_double2(fabs(a.x), fabs(a.y));
252 }
253 
254 
255 EIGEN_DEVICE_FUNC inline void
256 ptranspose(PacketBlock<float4,4>& kernel) {
257  double tmp = kernel.packet[0].y;
258  kernel.packet[0].y = kernel.packet[1].x;
259  kernel.packet[1].x = tmp;
260 
261  tmp = kernel.packet[0].z;
262  kernel.packet[0].z = kernel.packet[2].x;
263  kernel.packet[2].x = tmp;
264 
265  tmp = kernel.packet[0].w;
266  kernel.packet[0].w = kernel.packet[3].x;
267  kernel.packet[3].x = tmp;
268 
269  tmp = kernel.packet[1].z;
270  kernel.packet[1].z = kernel.packet[2].y;
271  kernel.packet[2].y = tmp;
272 
273  tmp = kernel.packet[1].w;
274  kernel.packet[1].w = kernel.packet[3].y;
275  kernel.packet[3].y = tmp;
276 
277  tmp = kernel.packet[2].w;
278  kernel.packet[2].w = kernel.packet[3].z;
279  kernel.packet[3].z = tmp;
280 }
281 
282 EIGEN_DEVICE_FUNC inline void
283 ptranspose(PacketBlock<double2,2>& kernel) {
284  double tmp = kernel.packet[0].y;
285  kernel.packet[0].y = kernel.packet[1].x;
286  kernel.packet[1].x = tmp;
287 }
288 
289 #endif
290 
291 } // end namespace internal
292 
293 } // end namespace Eigen
294 
295 
296 #endif // EIGEN_PACKET_MATH_CUDA_H
Definition: LDLT.h:16
Definition: Constants.h:222
Definition: Eigen_Colamd.h:54