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 template<> struct packet_traits<float> : default_packet_traits
25 {
26   typedef float4 type;
27   typedef float4 half;
28   enum {
29     Vectorizable = 1,
30     AlignedOnScalar = 1,
31     size=4,
32     HasHalfPacket = 0,
33 
34     HasDiv  = 1,
35     HasSin  = 0,
36     HasCos  = 0,
37     HasLog  = 1,
38     HasExp  = 1,
39     HasSqrt = 1,
40     HasRsqrt = 1,
41     HasLGamma = 1,
42     HasDiGamma = 1,
43     HasZeta = 1,
44     HasPolygamma = 1,
45     HasErf = 1,
46     HasErfc = 1,
47     HasIGamma = 1,
48     HasIGammac = 1,
49     HasBetaInc = 1,
50 
51     HasBlend = 0,
52   };
53 };
54 
55 template<> struct packet_traits<double> : default_packet_traits
56 {
57   typedef double2 type;
58   typedef double2 half;
59   enum {
60     Vectorizable = 1,
61     AlignedOnScalar = 1,
62     size=2,
63     HasHalfPacket = 0,
64 
65     HasDiv  = 1,
66     HasLog  = 1,
67     HasExp  = 1,
68     HasSqrt = 1,
69     HasRsqrt = 1,
70     HasLGamma = 1,
71     HasDiGamma = 1,
72     HasZeta = 1,
73     HasPolygamma = 1,
74     HasErf = 1,
75     HasErfc = 1,
76     HasIGamma = 1,
77     HasIGammac = 1,
78     HasBetaInc = 1,
79 
80     HasBlend = 0,
81   };
82 };
83 
84 
85 template<> struct unpacket_traits<float4>  { typedef float  type; enum {size=4, alignment=Aligned16}; typedef float4 half; };
86 template<> struct unpacket_traits<double2> { typedef double type; enum {size=2, alignment=Aligned16}; typedef double2 half; };
87 
88 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pset1<float4>(const float&  from) {
89   return make_float4(from, from, from, from);
90 }
91 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 pset1<double2>(const double& from) {
92   return make_double2(from, from);
93 }
94 
95 
96 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 plset<float4>(const float& a) {
97   return make_float4(a, a+1, a+2, a+3);
98 }
99 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 plset<double2>(const double& a) {
100   return make_double2(a, a+1);
101 }
102 
103 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 padd<float4>(const float4& a, const float4& b) {
104   return make_float4(a.x+b.x, a.y+b.y, a.z+b.z, a.w+b.w);
105 }
106 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 padd<double2>(const double2& a, const double2& b) {
107   return make_double2(a.x+b.x, a.y+b.y);
108 }
109 
110 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 psub<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 psub<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 pnegate(const float4& a) {
118   return make_float4(-a.x, -a.y, -a.z, -a.w);
119 }
120 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 pnegate(const double2& a) {
121   return make_double2(-a.x, -a.y);
122 }
123 
124 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pconj(const float4& a) { return a; }
125 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 pconj(const double2& a) { return a; }
126 
127 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pmul<float4>(const float4& a, const float4& b) {
128   return make_float4(a.x*b.x, a.y*b.y, a.z*b.z, a.w*b.w);
129 }
130 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 pmul<double2>(const double2& a, const double2& b) {
131   return make_double2(a.x*b.x, a.y*b.y);
132 }
133 
134 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pdiv<float4>(const float4& a, const float4& b) {
135   return make_float4(a.x/b.x, a.y/b.y, a.z/b.z, a.w/b.w);
136 }
137 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 pdiv<double2>(const double2& a, const double2& b) {
138   return make_double2(a.x/b.x, a.y/b.y);
139 }
140 
141 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pmin<float4>(const float4& a, const float4& b) {
142   return make_float4(fminf(a.x, b.x), fminf(a.y, b.y), fminf(a.z, b.z), fminf(a.w, b.w));
143 }
144 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 pmin<double2>(const double2& a, const double2& b) {
145   return make_double2(fmin(a.x, b.x), fmin(a.y, b.y));
146 }
147 
148 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pmax<float4>(const float4& a, const float4& b) {
149   return make_float4(fmaxf(a.x, b.x), fmaxf(a.y, b.y), fmaxf(a.z, b.z), fmaxf(a.w, b.w));
150 }
151 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 pmax<double2>(const double2& a, const double2& b) {
152   return make_double2(fmax(a.x, b.x), fmax(a.y, b.y));
153 }
154 
155 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pload<float4>(const float* from) {
156   return *reinterpret_cast<const float4*>(from);
157 }
158 
159 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 pload<double2>(const double* from) {
160   return *reinterpret_cast<const double2*>(from);
161 }
162 
163 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 ploadu<float4>(const float* from) {
164   return make_float4(from[0], from[1], from[2], from[3]);
165 }
166 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double2 ploadu<double2>(const double* from) {
167   return make_double2(from[0], from[1]);
168 }
169 
170 template<> EIGEN_STRONG_INLINE float4 ploaddup<float4>(const float*   from) {
171   return make_float4(from[0], from[0], from[1], from[1]);
172 }
173 template<> EIGEN_STRONG_INLINE double2 ploaddup<double2>(const double*  from) {
174   return make_double2(from[0], from[0]);
175 }
176 
177 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void pstore<float>(float*   to, const float4& from) {
178   *reinterpret_cast<float4*>(to) = from;
179 }
180 
181 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void pstore<double>(double* to, const double2& from) {
182   *reinterpret_cast<double2*>(to) = from;
183 }
184 
185 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void pstoreu<float>(float*  to, const float4& from) {
186   to[0] = from.x;
187   to[1] = from.y;
188   to[2] = from.z;
189   to[3] = from.w;
190 }
191 
192 template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void pstoreu<double>(double* to, const double2& from) {
193   to[0] = from.x;
194   to[1] = from.y;
195 }
196 
197 template<>
198 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float4 ploadt_ro<float4, Aligned>(const float* from) {
199 #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 350
200   return __ldg((const float4*)from);
201 #else
202   return make_float4(from[0], from[1], from[2], from[3]);
203 #endif
204 }
205 template<>
206 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double2 ploadt_ro<double2, Aligned>(const double* from) {
207 #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 350
208   return __ldg((const double2*)from);
209 #else
210   return make_double2(from[0], from[1]);
211 #endif
212 }
213 
214 template<>
215 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float4 ploadt_ro<float4, Unaligned>(const float* from) {
216 #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 350
217   return make_float4(__ldg(from+0), __ldg(from+1), __ldg(from+2), __ldg(from+3));
218 #else
219   return make_float4(from[0], from[1], from[2], from[3]);
220 #endif
221 }
222 template<>
223 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double2 ploadt_ro<double2, Unaligned>(const double* from) {
224 #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 350
225   return make_double2(__ldg(from+0), __ldg(from+1));
226 #else
227   return make_double2(from[0], from[1]);
228 #endif
229 }
230 
231 template<> EIGEN_DEVICE_FUNC inline float4 pgather<float, float4>(const float* from, Index stride) {
232   return make_float4(from[0*stride], from[1*stride], from[2*stride], from[3*stride]);
233 }
234 
235 template<> EIGEN_DEVICE_FUNC inline double2 pgather<double, double2>(const double* from, Index stride) {
236   return make_double2(from[0*stride], from[1*stride]);
237 }
238 
239 template<> EIGEN_DEVICE_FUNC inline void pscatter<float, float4>(float* to, const float4& from, Index stride) {
240   to[stride*0] = from.x;
241   to[stride*1] = from.y;
242   to[stride*2] = from.z;
243   to[stride*3] = from.w;
244 }
245 template<> EIGEN_DEVICE_FUNC inline void pscatter<double, double2>(double* to, const double2& from, Index stride) {
246   to[stride*0] = from.x;
247   to[stride*1] = from.y;
248 }
249 
250 template<> EIGEN_DEVICE_FUNC inline float  pfirst<float4>(const float4& a) {
251   return a.x;
252 }
253 template<> EIGEN_DEVICE_FUNC inline double pfirst<double2>(const double2& a) {
254   return a.x;
255 }
256 
257 template<> EIGEN_DEVICE_FUNC inline float  predux<float4>(const float4& a) {
258   return a.x + a.y + a.z + a.w;
259 }
260 template<> EIGEN_DEVICE_FUNC inline double predux<double2>(const double2& a) {
261   return a.x + a.y;
262 }
263 
264 template<> EIGEN_DEVICE_FUNC inline float  predux_max<float4>(const float4& a) {
265   return fmaxf(fmaxf(a.x, a.y), fmaxf(a.z, a.w));
266 }
267 template<> EIGEN_DEVICE_FUNC inline double predux_max<double2>(const double2& a) {
268   return fmax(a.x, a.y);
269 }
270 
271 template<> EIGEN_DEVICE_FUNC inline float  predux_min<float4>(const float4& a) {
272   return fminf(fminf(a.x, a.y), fminf(a.z, a.w));
273 }
274 template<> EIGEN_DEVICE_FUNC inline double predux_min<double2>(const double2& a) {
275   return fmin(a.x, a.y);
276 }
277 
278 template<> EIGEN_DEVICE_FUNC inline float  predux_mul<float4>(const float4& a) {
279   return a.x * a.y * a.z * a.w;
280 }
281 template<> EIGEN_DEVICE_FUNC inline double predux_mul<double2>(const double2& a) {
282   return a.x * a.y;
283 }
284 
285 template<> EIGEN_DEVICE_FUNC inline float4  pabs<float4>(const float4& a) {
286   return make_float4(fabsf(a.x), fabsf(a.y), fabsf(a.z), fabsf(a.w));
287 }
288 template<> EIGEN_DEVICE_FUNC inline double2 pabs<double2>(const double2& a) {
289   return make_double2(fabs(a.x), fabs(a.y));
290 }
291 
292 EIGEN_DEVICE_FUNC inline void
293 ptranspose(PacketBlock<float4,4>& kernel) {
294   float tmp = kernel.packet[0].y;
295   kernel.packet[0].y = kernel.packet[1].x;
296   kernel.packet[1].x = tmp;
297 
298   tmp = kernel.packet[0].z;
299   kernel.packet[0].z = kernel.packet[2].x;
300   kernel.packet[2].x = tmp;
301 
302   tmp = kernel.packet[0].w;
303   kernel.packet[0].w = kernel.packet[3].x;
304   kernel.packet[3].x = tmp;
305 
306   tmp = kernel.packet[1].z;
307   kernel.packet[1].z = kernel.packet[2].y;
308   kernel.packet[2].y = tmp;
309 
310   tmp = kernel.packet[1].w;
311   kernel.packet[1].w = kernel.packet[3].y;
312   kernel.packet[3].y = tmp;
313 
314   tmp = kernel.packet[2].w;
315   kernel.packet[2].w = kernel.packet[3].z;
316   kernel.packet[3].z = tmp;
317 }
318 
319 EIGEN_DEVICE_FUNC inline void
320 ptranspose(PacketBlock<double2,2>& kernel) {
321   double tmp = kernel.packet[0].y;
322   kernel.packet[0].y = kernel.packet[1].x;
323   kernel.packet[1].x = tmp;
324 }
325 
326 #endif
327 
328 } // end namespace internal
329 
330 } // end namespace Eigen
331 
332 
333 #endif // EIGEN_PACKET_MATH_CUDA_H
334