1 #include <unittest/unittest.h>
2 #include <thrust/host_vector.h>
3 #include <thrust/complex.h>
4 #include <thrust/transform.h>
5 #include <iostream>
6
7 #if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
8 #include <backend/cuda/testframework.h>
9 #endif
10
11 struct basic_arithmetic_functor
12 {
13 template<typename T>
14 __host__ __device__
operator ()basic_arithmetic_functor15 thrust::complex<T> operator()(const thrust::complex<T> &x,
16 const thrust::complex<T> &y)
17 {
18 // exercise unary and binary arithmetic operators
19 // Should return approximately 1
20 return (+x + +y) + (x * y) / (y * x) + (-y + -x);
21 } // end operator()()
22 }; // end make_pair_functor
23
24 struct complex_plane_functor
25 {
26 template<typename T>
27 __host__ __device__
operator ()complex_plane_functor28 thrust::complex<T> operator()(const thrust::complex<T> &x)
29 {
30 // Should return a proximately 1
31 return thrust::proj( (thrust::polar(abs(x),arg(x)) * conj(x))/norm(x));
32 } // end operator()()
33 }; // end make_pair_functor
34
35 struct pow_functor
36 {
37 template<typename T>
38 __host__ __device__
operator ()pow_functor39 thrust::complex<T> operator()(const thrust::complex<T> &x,
40 const thrust::complex<T> &y)
41 {
42 // exercise power functions
43 return pow(x,y);
44 } // end operator()()
45 }; // end make_pair_functor
46
47 struct sqrt_functor
48 {
49 template<typename T>
50 __host__ __device__
operator ()sqrt_functor51 thrust::complex<T> operator()(const thrust::complex<T> &x)
52 {
53 // exercise power functions
54 return sqrt(x);
55 } // end operator()()
56 }; // end make_pair_functor
57
58 struct log_functor
59 {
60 template<typename T>
61 __host__ __device__
operator ()log_functor62 thrust::complex<T> operator()(const thrust::complex<T> &x)
63 {
64 return log(x);
65 } // end operator()()
66 }; // end make_pair_functor
67
68 struct exp_functor
69 {
70 template<typename T>
71 __host__ __device__
operator ()exp_functor72 thrust::complex<T> operator()(const thrust::complex<T> &x)
73 {
74 return exp(x);
75 } // end operator()()
76 }; // end make_pair_functor
77
78 struct log10_functor
79 {
80 template<typename T>
81 __host__ __device__
operator ()log10_functor82 thrust::complex<T> operator()(const thrust::complex<T> &x)
83 {
84 return log10(x);
85 } // end operator()()
86 }; // end make_pair_functor
87
88
89 struct cos_functor
90 {
91 template<typename T>
92 __host__ __device__
operator ()cos_functor93 thrust::complex<T> operator()(const thrust::complex<T> &x)
94 {
95 return cos(x);
96 }
97 };
98
99 struct sin_functor
100 {
101 template<typename T>
102 __host__ __device__
operator ()sin_functor103 thrust::complex<T> operator()(const thrust::complex<T> &x)
104 {
105 return sin(x);
106 }
107 };
108
109 struct tan_functor
110 {
111 template<typename T>
112 __host__ __device__
operator ()tan_functor113 thrust::complex<T> operator()(const thrust::complex<T> &x)
114 {
115 return tan(x);
116 }
117 };
118
119
120
121 struct cosh_functor
122 {
123 template<typename T>
124 __host__ __device__
operator ()cosh_functor125 thrust::complex<T> operator()(const thrust::complex<T> &x)
126 {
127 return cosh(x);
128 }
129 };
130
131 struct sinh_functor
132 {
133 template<typename T>
134 __host__ __device__
operator ()sinh_functor135 thrust::complex<T> operator()(const thrust::complex<T> &x)
136 {
137 return sinh(x);
138 }
139 };
140
141 struct tanh_functor
142 {
143 template<typename T>
144 __host__ __device__
operator ()tanh_functor145 thrust::complex<T> operator()(const thrust::complex<T> &x)
146 {
147 return tanh(x);
148 }
149 };
150
151
152 struct acos_functor
153 {
154 template<typename T>
155 __host__ __device__
operator ()acos_functor156 thrust::complex<T> operator()(const thrust::complex<T> &x)
157 {
158 return acos(x);
159 }
160 };
161
162 struct asin_functor
163 {
164 template<typename T>
165 __host__ __device__
operator ()asin_functor166 thrust::complex<T> operator()(const thrust::complex<T> &x)
167 {
168 return asin(x);
169 }
170 };
171
172 struct atan_functor
173 {
174 template<typename T>
175 __host__ __device__
operator ()atan_functor176 thrust::complex<T> operator()(const thrust::complex<T> &x)
177 {
178 return atan(x);
179 }
180 };
181
182
183 struct acosh_functor
184 {
185 template<typename T>
186 __host__ __device__
operator ()acosh_functor187 thrust::complex<T> operator()(const thrust::complex<T> &x)
188 {
189 return acosh(x);
190 }
191 };
192
193 struct asinh_functor
194 {
195 template<typename T>
196 __host__ __device__
operator ()asinh_functor197 thrust::complex<T> operator()(const thrust::complex<T> &x)
198 {
199 return asinh(x);
200 }
201 };
202
203 struct atanh_functor
204 {
205 template<typename T>
206 __host__ __device__
operator ()atanh_functor207 thrust::complex<T> operator()(const thrust::complex<T> &x)
208 {
209 return atanh(x);
210 }
211 };
212
213
214 template <typename T>
random_complex_samples(size_t n)215 thrust::host_vector<thrust::complex<T> > random_complex_samples(size_t n){
216 thrust::host_vector<T> real = unittest::random_samples<T>(2*n);
217 thrust::host_vector<thrust::complex<T> > h_p1(n);
218 for(size_t i = 0; i<n; i++){
219 h_p1[i].real(real[i]);
220 h_p1[i].imag(real[2*i]);
221 }
222 return h_p1;
223 }
224
225 template <typename T>
226 struct TestComplexArithmeticTransform
227 {
operator ()TestComplexArithmeticTransform228 void operator()(const size_t n)
229 {
230 typedef thrust::complex<T> type;
231 thrust::host_vector<type> h_p1 = random_complex_samples<T>(n);
232 thrust::host_vector<type> h_p2 = random_complex_samples<T>(n);
233 thrust::host_vector<type> h_result(n);
234
235 thrust::device_vector<type> d_p1 = h_p1;
236 thrust::device_vector<type> d_p2 = h_p2;
237 thrust::device_vector<type> d_result(n);
238 #if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
239 if(const CUDATestDriver *driver = dynamic_cast<const CUDATestDriver*>(&UnitTestDriver::s_driver()))
240 {
241 if(sizeof(T) == sizeof(double) && driver->current_device_architecture() < 200)
242 {
243 KNOWN_FAILURE;
244 } // end if
245 } // end if
246 #endif
247
248 thrust::transform(h_p1.begin(), h_p1.end(), h_p2.begin(), h_result.begin(), basic_arithmetic_functor());
249 thrust::transform(d_p1.begin(), d_p1.end(), d_p2.begin(), d_result.begin(), basic_arithmetic_functor());
250 ASSERT_ALMOST_EQUAL(h_result, d_result);
251 }
252 };
253 VariableUnitTest<TestComplexArithmeticTransform, FloatingPointTypes> TestComplexArithmeticTransformInstance;
254
255 template <typename T>
256 struct TestComplexPlaneTransform
257 {
operator ()TestComplexPlaneTransform258 void operator()(const size_t n)
259 {
260 typedef thrust::complex<T> type;
261 thrust::host_vector<type> h_p1 = random_complex_samples<T>(n);
262 thrust::host_vector<type> h_result(n);
263
264 thrust::device_vector<type> d_p1 = h_p1;
265 thrust::device_vector<type> d_result(n);
266
267 #if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
268 if(const CUDATestDriver *driver = dynamic_cast<const CUDATestDriver*>(&UnitTestDriver::s_driver()))
269 {
270 if(sizeof(T) == sizeof(double) && driver->current_device_architecture() < 200)
271 {
272 KNOWN_FAILURE;
273 } // end if
274 } // end if
275 #endif
276
277 thrust::transform(h_p1.begin(), h_p1.end(), h_result.begin(), complex_plane_functor());
278 thrust::transform(d_p1.begin(), d_p1.end(), d_result.begin(), complex_plane_functor());
279 ASSERT_ALMOST_EQUAL(h_result, d_result);
280 }
281 };
282 VariableUnitTest<TestComplexPlaneTransform, FloatingPointTypes> TestComplexPlaneTransformInstance;
283
284
285 template <typename T>
286 struct TestComplexPowerTransform
287 {
operator ()TestComplexPowerTransform288 void operator()(const size_t n)
289 {
290 typedef thrust::complex<T> type;
291 thrust::host_vector<type> h_p1 = random_complex_samples<T>(n);
292 thrust::host_vector<type> h_p2 = random_complex_samples<T>(n);
293 thrust::host_vector<type> h_result(n);
294
295 thrust::device_vector<type> d_p1 = h_p1;
296 thrust::device_vector<type> d_p2 = h_p2;
297 thrust::device_vector<type> d_result(n);
298
299 #if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
300 if(const CUDATestDriver *driver = dynamic_cast<const CUDATestDriver*>(&UnitTestDriver::s_driver()))
301 {
302 if(sizeof(T) == sizeof(double) && driver->current_device_architecture() < 200)
303 {
304 KNOWN_FAILURE;
305 } // end if
306 } // end if
307 #endif
308
309 thrust::transform(h_p1.begin(), h_p1.end(), h_p2.begin(), h_result.begin(), pow_functor());
310 thrust::transform(d_p1.begin(), d_p1.end(), d_p2.begin(), d_result.begin(), pow_functor());
311 // pow can be very innacurate there's no point trying to check for equality
312 // Currently just checking for compilation
313 // ASSERT_ALMOST_EQUAL(h_result, d_result);
314
315 thrust::transform(h_p1.begin(), h_p1.end(), h_result.begin(), sqrt_functor());
316 thrust::transform(d_p1.begin(), d_p1.end(), d_result.begin(), sqrt_functor());
317 ASSERT_ALMOST_EQUAL(h_result, d_result);
318 }
319 };
320 VariableUnitTest<TestComplexPowerTransform, FloatingPointTypes> TestComplexPowerTransformInstance;
321
322 template <typename T>
323 struct TestComplexExponentialTransform
324 {
operator ()TestComplexExponentialTransform325 void operator()(const size_t n)
326 {
327 typedef thrust::complex<T> type;
328 thrust::host_vector<type> h_p1 = random_complex_samples<T>(n);
329 thrust::host_vector<type> h_result(n);
330
331 thrust::device_vector<type> d_p1 = h_p1;
332 thrust::device_vector<type> d_result(n);
333
334 #if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
335 if(const CUDATestDriver *driver = dynamic_cast<const CUDATestDriver*>(&UnitTestDriver::s_driver()))
336 {
337 if(sizeof(T) == sizeof(double) && driver->current_device_architecture() < 200)
338 {
339 KNOWN_FAILURE;
340 } // end if
341 } // end if
342 #endif
343
344 thrust::transform(h_p1.begin(), h_p1.end(), h_result.begin(), exp_functor());
345 thrust::transform(d_p1.begin(), d_p1.end(), d_result.begin(), exp_functor());
346 ASSERT_ALMOST_EQUAL(h_result, d_result);
347
348 thrust::transform(h_p1.begin(), h_p1.end(), h_result.begin(), log_functor());
349 thrust::transform(d_p1.begin(), d_p1.end(), d_result.begin(), log_functor());
350 ASSERT_ALMOST_EQUAL(h_result, d_result);
351
352 thrust::transform(h_p1.begin(), h_p1.end(), h_result.begin(), log10_functor());
353 thrust::transform(d_p1.begin(), d_p1.end(), d_result.begin(), log10_functor());
354 ASSERT_ALMOST_EQUAL(h_result, d_result);
355 }
356 };
357 VariableUnitTest<TestComplexExponentialTransform, FloatingPointTypes> TestComplexExponentialTransformInstance;
358
359 template <typename T>
360 struct TestComplexTrigonometricTransform
361 {
operator ()TestComplexTrigonometricTransform362 void operator()(const size_t n)
363 {
364 typedef thrust::complex<T> type;
365 thrust::host_vector<type> h_p1 = random_complex_samples<T>(n);
366 thrust::host_vector<type> h_result(n);
367
368 thrust::device_vector<type> d_p1 = h_p1;
369 thrust::device_vector<type> d_result(n);
370
371 #if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
372 if(const CUDATestDriver *driver = dynamic_cast<const CUDATestDriver*>(&UnitTestDriver::s_driver()))
373 {
374 if(sizeof(T) == sizeof(double) && driver->current_device_architecture() < 200)
375 {
376 KNOWN_FAILURE;
377 } // end if
378 } // end if
379 #endif
380
381 thrust::transform(h_p1.begin(), h_p1.end(), h_result.begin(), sin_functor());
382 thrust::transform(d_p1.begin(), d_p1.end(), d_result.begin(), sin_functor());
383 ASSERT_ALMOST_EQUAL(h_result, d_result);
384
385 thrust::transform(h_p1.begin(), h_p1.end(), h_result.begin(), cos_functor());
386 thrust::transform(d_p1.begin(), d_p1.end(), d_result.begin(), cos_functor());
387 ASSERT_ALMOST_EQUAL(h_result, d_result);
388
389 thrust::transform(h_p1.begin(), h_p1.end(), h_result.begin(), tan_functor());
390 thrust::transform(d_p1.begin(), d_p1.end(), d_result.begin(), tan_functor());
391 ASSERT_ALMOST_EQUAL(h_result, d_result);
392
393
394 thrust::transform(h_p1.begin(), h_p1.end(), h_result.begin(), sinh_functor());
395 thrust::transform(d_p1.begin(), d_p1.end(), d_result.begin(), sinh_functor());
396 ASSERT_ALMOST_EQUAL(h_result, d_result);
397
398 thrust::transform(h_p1.begin(), h_p1.end(), h_result.begin(), cosh_functor());
399 thrust::transform(d_p1.begin(), d_p1.end(), d_result.begin(), cosh_functor());
400 ASSERT_ALMOST_EQUAL(h_result, d_result);
401
402 thrust::transform(h_p1.begin(), h_p1.end(), h_result.begin(), tanh_functor());
403 thrust::transform(d_p1.begin(), d_p1.end(), d_result.begin(), tanh_functor());
404 ASSERT_ALMOST_EQUAL(h_result, d_result);
405
406
407
408 thrust::transform(h_p1.begin(), h_p1.end(), h_result.begin(), asin_functor());
409 thrust::transform(d_p1.begin(), d_p1.end(), d_result.begin(), asin_functor());
410 ASSERT_ALMOST_EQUAL(h_result, d_result);
411
412 thrust::transform(h_p1.begin(), h_p1.end(), h_result.begin(), acos_functor());
413 thrust::transform(d_p1.begin(), d_p1.end(), d_result.begin(), acos_functor());
414 ASSERT_ALMOST_EQUAL(h_result, d_result);
415
416 thrust::transform(h_p1.begin(), h_p1.end(), h_result.begin(), atan_functor());
417 thrust::transform(d_p1.begin(), d_p1.end(), d_result.begin(), atan_functor());
418 ASSERT_ALMOST_EQUAL(h_result, d_result);
419
420
421 thrust::transform(h_p1.begin(), h_p1.end(), h_result.begin(), asinh_functor());
422 thrust::transform(d_p1.begin(), d_p1.end(), d_result.begin(), asinh_functor());
423 ASSERT_ALMOST_EQUAL(h_result, d_result);
424
425 thrust::transform(h_p1.begin(), h_p1.end(), h_result.begin(), acosh_functor());
426 thrust::transform(d_p1.begin(), d_p1.end(), d_result.begin(), acosh_functor());
427 ASSERT_ALMOST_EQUAL(h_result, d_result);
428
429 thrust::transform(h_p1.begin(), h_p1.end(), h_result.begin(), atanh_functor());
430 thrust::transform(d_p1.begin(), d_p1.end(), d_result.begin(), atanh_functor());
431 ASSERT_ALMOST_EQUAL(h_result, d_result);
432
433 }
434 };
435 VariableUnitTest<TestComplexTrigonometricTransform, FloatingPointTypes> TestComplexTrigonometricTransformInstance;
436
437