Fermat
numbers.h
1 /*
2  * CUGAR : Cuda Graphics Accelerator
3  *
4  * Copyright (c) 2011-2019, NVIDIA CORPORATION. All rights reserved.
5  *
6  * Redistribution and use in source and binary forms, with or without
7  * modification, are permitted provided that the following conditions are met:
8  * * Redistributions of source code must retain the above copyright
9  * notice, this list of conditions and the following disclaimer.
10  * * Redistributions in binary form must reproduce the above copyright
11  * notice, this list of conditions and the following disclaimer in the
12  * documentation and/or other materials provided with the distribution.
13  * * Neither the name of the NVIDIA CORPORATION nor the
14  * names of its contributors may be used to endorse or promote products
15  * derived from this software without specific prior written permission.
16  *
17  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
18  * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
19  * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
20  * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
21  * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
22  * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
23  * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
24  * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
25  * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
26  * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
27  */
28 
29 #pragma once
30 
31 #include <cmath>
32 #include <limits>
33 #include <cugar/basic/types.h>
34 #include <vector_types.h>
35 #include <vector_functions.h>
36 #include <cuda_fp16.h>
37 
38 namespace cugar {
39 
40 #if defined(WIN32)
41 #include <float.h>
42 
43 #ifndef M_PI
44 #define M_PI 3.14159265358979323846
45 #endif
46 
47 #ifndef M_PIf
48 #define M_PIf 3.14159265358979323846f
49 #endif
50 
51 #ifndef M_TWO_PI
52 #define M_TWO_PI 6.28318530717958647693
53 #endif
54 
55 #ifndef M_TWO_PIf
56 #define M_TWO_PIf 6.28318530717958647693f
57 #endif
58 
59 CUGAR_HOST_DEVICE inline bool is_finite(const double x)
60 {
61  #if defined(CUGAR_DEVICE_COMPILATION)
62  return isfinite(x) != 0;
63  #else
64  return _finite(x) != 0;
65  #endif
66 }
67 CUGAR_HOST_DEVICE
68 inline bool is_nan(const double x)
69 {
70  #if defined(CUGAR_DEVICE_COMPILATION)
71  return isnan(x) != 0;
72  #else
73  return _isnan(x) != 0;
74  #endif
75 }
76 CUGAR_HOST_DEVICE inline bool is_finite(const float x)
77 {
78  #if defined(CUGAR_DEVICE_COMPILATION)
79  return isfinite(x) != 0;
80  #else
81  return _finite(x) != 0;
82  #endif
83 }
84 CUGAR_HOST_DEVICE inline bool is_nan(const float x)
85 {
86  #if defined(CUGAR_DEVICE_COMPILATION)
87  return isnan(x) != 0;
88  #else
89  return _isnan(x) != 0;
90  #endif
91 }
92 
93 #else
94 
95 #ifndef M_PI
96 #define M_PI 3.14159265358979323846
97 #endif
98 
99 #ifndef M_TWO_PI
100 #define M_TWO_PI 6.28318530717958647693
101 #endif
102 
103 #ifndef M_PIf
104 #define M_PIf 3.14159265358979323846f
105 #endif
106 
107 #ifndef M_TWO_PIf
108 #define M_TWO_PIf 6.28318530717958647693f
109 #endif
110 
111 #endif
112 
113 #ifdef __CUDACC__
114 
115 CUGAR_FORCEINLINE __device__ uint32 warp_tid() { return threadIdx.x & 31; }
116 CUGAR_FORCEINLINE __device__ uint32 warp_id() { return threadIdx.x >> 5; }
117 
118 #endif
119 
120 CUGAR_HOST_DEVICE CUGAR_FORCEINLINE
121 float float_infinity() { return cugar::binary_cast<float>(0x7f800000u); }
122 
123 CUGAR_HOST_DEVICE CUGAR_FORCEINLINE
124 double double_infinity() { return cugar::binary_cast<double>(0x7ff0000000000000ULL ); }
125 
134 
137 
142 
146 template <uint32 N>
147 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE
148 uint32 lo_bits() { return (1u << N) - 1u; }
149 
153 template <uint32 N>
154 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE
155 uint32 hi_bits() { return ~lo_bits<N>(); }
156 
160 template <typename Iterator, typename T>
161 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE
162 uint32 count_occurrences(const Iterator begin, uint32 size, const T val, const uint32 max_occ = uint32(-1))
163 {
164  uint32 occ = 0u;
165  for (uint32 i = 0; i < size; ++i)
166  {
167  if (begin[i] == val)
168  {
169  if (++occ >= max_occ)
170  return occ;
171  }
172  }
173  return occ;
174 }
175 
179 template<typename L, typename R>
180 inline CUGAR_HOST_DEVICE L divide_ri(const L x, const R y)
181 {
182  return L( (x + (y - 1)) / y );
183 }
184 
188 template<typename L, typename R>
189 inline CUGAR_HOST_DEVICE L divide_rz(const L x, const R y)
190 {
191  return L( x / y );
192 }
193 
197 template<typename L, typename R>
198 inline CUGAR_HOST_DEVICE L round_i(const L x, const R y){ return L( y * divide_ri(x, y) ); }
199 
203 template<typename L, typename R>
204 inline CUGAR_HOST_DEVICE L round_z(const L x, const R y){ return L( y * divide_rz(x, y) ); }
205 
209 template<typename L, typename R>
210 inline CUGAR_HOST_DEVICE L round(const L x, const R y)
211 {
212  const L r = round_z( x, y );
213  return R((x - r)*2) > y ? r+L(1) : r;
214 }
215 
218 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE uint8 comp(const uchar2 a, const char c)
219 {
220  return (c == 0 ? a.x : a.y);
221 }
224 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE char comp(const char2 a, const char c)
225 {
226  return (c == 0 ? a.x : a.y);
227 }
228 
231 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE uint8 comp(const uchar4 a, const char c)
232 {
233  return c <= 1 ?
234  (c == 0 ? a.x : a.y) :
235  (c == 2 ? a.z : a.w);
236 }
239 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE char comp(const char4 a, const char c)
240 {
241  return c <= 1 ?
242  (c == 0 ? a.x : a.y) :
243  (c == 2 ? a.z : a.w);
244 }
245 
248 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE uint32 comp(const uint2 a, const uint32 c)
249 {
250  return (c == 0 ? a.x : a.y);
251 }
254 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE void set(uint2& a, const uint32 c, const uint32 v)
255 {
256  if (c == 0) a.x = v;
257  else a.y = v;
258 }
261 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE uint64 comp(const ulonglong2 a, const uint32 c)
262 {
263  return (c == 0 ? a.x : a.y);
264 }
267 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE void set(ulonglong2& a, const uint32 c, const uint64 v)
268 {
269  if (c == 0) a.x = v;
270  else a.y = v;
271 }
274 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE int32 comp(const int2 a, const uint32 c)
275 {
276  return (c == 0 ? a.x : a.y);
277 }
280 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE uint32 comp(const uint4 a, const uint32 c)
281 {
282  return c <= 1 ?
283  (c == 0 ? a.x : a.y) :
284  (c == 2 ? a.z : a.w);
285 }
288 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE void set(uint4& a, const uint32 c, const uint32 v)
289 {
290  if (c == 0) a.x = v;
291  else if (c == 1) a.y = v;
292  else if (c == 2) a.z = v;
293  else a.w = v;
294 }
297 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE int32 comp(const int4 a, const uint32 c)
298 {
299  return c <= 1 ?
300  (c == 0 ? a.x : a.y) :
301  (c == 2 ? a.z : a.w);
302 }
303 
306 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE uint16 comp(const ushort4 a, const uint32 c)
307 {
308  return c <= 1 ?
309  (c == 0 ? a.x : a.y) :
310  (c == 2 ? a.z : a.w);
311 }
312 
315 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE uint64 comp(const ulonglong4 a, const uint32 c)
316 {
317  return c <= 1 ?
318  (c == 0 ? a.x : a.y) :
319  (c == 2 ? a.z : a.w);
320 }
323 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE void set(ulonglong4& a, const uint32 c, const uint64 v)
324 {
325  if (c == 0) a.x = v;
326  else if (c == 1) a.y = v;
327  else if (c == 2) a.z = v;
328  else a.w = v;
329 }
330 
334 
335 typedef uchar2 uint8_2;
336 typedef uchar3 uint8_3;
337 typedef uchar4 uint8_4;
338 
339 typedef char2 int8_2;
340 typedef char3 int8_3;
341 typedef char4 int8_4;
342 
343 typedef ushort2 uint16_2;
344 typedef ushort3 uint16_3;
345 typedef ushort4 uint16_4;
346 
347 typedef short2 int16_2;
348 typedef short3 int16_3;
349 typedef short4 int16_4;
350 
351 typedef uint2 uint32_2;
352 typedef uint3 uint32_3;
353 typedef uint4 uint32_4;
354 
355 typedef int2 int32_2;
356 typedef int3 int32_3;
357 typedef int4 int32_4;
358 
359 typedef ulonglong2 uint64_2;
360 typedef ulonglong3 uint64_3;
361 typedef ulonglong4 uint64_4;
362 
363 typedef longlong2 int64_2;
364 typedef longlong3 int64_3;
365 typedef longlong4 int64_4;
366 
367 template <typename T, uint32 DIM>
368 struct vector_type {};
369 
370 template <typename T>
371 struct vector1_storage
372 {
373  CUGAR_FORCEINLINE CUGAR_HOST_DEVICE vector1_storage() {}
374  CUGAR_FORCEINLINE CUGAR_HOST_DEVICE vector1_storage(T _x) : x(_x) {}
375 
376  T x;
377 };
378 template <typename T>
379 struct vector2_storage
380 {
381  CUGAR_FORCEINLINE CUGAR_HOST_DEVICE vector2_storage() {}
382  CUGAR_FORCEINLINE CUGAR_HOST_DEVICE vector2_storage(T _x, T _y) : x(_x), y(_y) {}
383 
384  T x, y;
385 };
386 template <typename T>
387 struct vector3_storage
388 {
389  CUGAR_FORCEINLINE CUGAR_HOST_DEVICE vector3_storage() {}
390  CUGAR_FORCEINLINE CUGAR_HOST_DEVICE vector3_storage(T _x, T _y, T _z) : x(_x), y(_y), z(_z) {}
391 
392  T x, y, z;
393 };
394 template <typename T>
395 struct vector4_storage
396 {
397  CUGAR_FORCEINLINE CUGAR_HOST_DEVICE vector4_storage() {}
398  CUGAR_FORCEINLINE CUGAR_HOST_DEVICE vector4_storage(T _x, T _y, T _z, T _w) : x(_x), y(_y), z(_z), w(_w) {}
399 
400  T x, y, z, w;
401 };
402 
403 template <typename T> struct vector_type<T,1> { typedef vector1_storage<T> type; static type make(const T i1) { return type(i1); } };
404 template <typename T> struct vector_type<T,2> { typedef vector2_storage<T> type; static type make(const T i1, const T i2) { return type(i1,i2); } };
405 template <typename T> struct vector_type<T,3> { typedef vector3_storage<T> type; static type make(const T i1, const T i2, const T i3) { return type(i1,i2,i3); } };
406 template <typename T> struct vector_type<T,4> { typedef vector4_storage<T> type; static type make(const T i1, const T i2, const T i3, const T i4) { return type(i1,i2,i3,4); } };
407 
408 template <> struct vector_type<char,1> { typedef char type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const char i1) { return i1; } };
409 template <> struct vector_type<char,2> { typedef char2 type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const char i1, const char i2) { return make_char2(i1,i2); } };
410 template <> struct vector_type<char,3> { typedef char3 type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const char i1, const char i2, const char i3) { return make_char3(i1,i2,i3); } };
411 template <> struct vector_type<char,4> { typedef char4 type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const char i1, const char i2, const char i3, const char i4) { return make_char4(i1,i2,i3,i4); } };
412 
413 template <> struct vector_type<unsigned char,1> { typedef unsigned char type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const unsigned char i1) { return i1; } };
414 template <> struct vector_type<unsigned char,2> { typedef uchar2 type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const unsigned char i1, const unsigned char i2) { return make_uchar2(i1,i2); } };
415 template <> struct vector_type<unsigned char,3> { typedef uchar3 type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const unsigned char i1, const unsigned char i2, const unsigned char i3) { return make_uchar3(i1,i2,i3); } };
416 template <> struct vector_type<unsigned char,4> { typedef uchar4 type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const unsigned char i1, const unsigned char i2, const unsigned char i3, const unsigned char i4) { return make_uchar4(i1,i2,i3,i4); } };
417 
418 template <> struct vector_type<short,1> { typedef short type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const short i1) { return i1; } };
419 template <> struct vector_type<short,2> { typedef short2 type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const short i1, const short i2) { return make_short2(i1,i2); } };
420 template <> struct vector_type<short,3> { typedef short3 type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const short i1, const short i2, const short i3) { return make_short3(i1,i2,i3); } };
421 template <> struct vector_type<short,4> { typedef short4 type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const short i1, const short i2, const short i3, const short i4) { return make_short4(i1,i2,i3,i4); } };
422 
423 template <> struct vector_type<unsigned short,1> { typedef unsigned short type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const unsigned short i1) { return i1; } };
424 template <> struct vector_type<unsigned short,2> { typedef ushort2 type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const unsigned short i1, const unsigned short i2) { return make_ushort2(i1,i2); } };
425 template <> struct vector_type<unsigned short,3> { typedef ushort3 type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const unsigned short i1, const unsigned short i2, const unsigned short i3) { return make_ushort3(i1,i2,i3); } };
426 template <> struct vector_type<unsigned short,4> { typedef ushort4 type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const unsigned short i1, const unsigned short i2, const unsigned short i3, const unsigned short i4) { return make_ushort4(i1,i2,i3,i4); } };
427 
428 template <> struct vector_type<int,1> { typedef int type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const int i1) { return i1; } };
429 template <> struct vector_type<int,2> { typedef int2 type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const int i1, const int i2) { return make_int2(i1,i2); } };
430 template <> struct vector_type<int,3> { typedef int3 type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const int i1, const int i2, const int i3) { return make_int3(i1,i2,i3); } };
431 template <> struct vector_type<int,4> { typedef int4 type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const int i1, const int i2, const int i3, const int i4) { return make_int4(i1,i2,i3,i4); } };
432 
433 template <> struct vector_type<unsigned int,1> { typedef unsigned int type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const unsigned int i1) { return i1; } };
434 template <> struct vector_type<unsigned int,2> { typedef uint2 type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const unsigned int i1, const unsigned int i2) { return make_uint2(i1,i2); } };
435 template <> struct vector_type<unsigned int,3> { typedef uint3 type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const unsigned int i1, const unsigned int i2, const unsigned int i3) { return make_uint3(i1,i2,i3); } };
436 template <> struct vector_type<unsigned int,4> { typedef uint4 type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const unsigned int i1, const unsigned int i2, const unsigned int i3, const unsigned int i4) { return make_uint4(i1,i2,i3,i4); } };
437 
438 template <> struct vector_type<int64,1> { typedef int64 type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const int64 i1) { return i1; } };
439 template <> struct vector_type<int64,2> { typedef int64_2 type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const int64 i1, const int64 i2) { int64_2 r; r.x = i1; r.y = i2; return r; } };
440 template <> struct vector_type<int64,3> { typedef int64_3 type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const int64 i1, const int64 i2, const int64 i3) { int64_3 r; r.x = i1; r.y = i2; r.z = i3; return r; } };
441 template <> struct vector_type<int64,4> { typedef int64_4 type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const int64 i1, const int64 i2, const int64 i3, const int64 i4) { int64_4 r; r.x = i1; r.y = i2; r.z = i3, r.w = i4; return r; } };
442 
443 template <> struct vector_type<uint64,1> { typedef uint64 type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const uint64 i1) { return i1; } };
444 template <> struct vector_type<uint64,2> { typedef uint64_2 type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const uint64 i1, const uint64 i2) { uint64_2 r; r.x = i1; r.y = i2; return r; } };
445 template <> struct vector_type<uint64,3> { typedef uint64_3 type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const uint64 i1, const uint64 i2, const uint64 i3) { uint64_3 r; r.x = i1; r.y = i2; r.z = i3; return r; } };
446 template <> struct vector_type<uint64,4> { typedef uint64_4 type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const uint64 i1, const uint64 i2, const uint64 i3, const uint64 i4) { uint64_4 r; r.x = i1; r.y = i2; r.z = i3, r.w = i4; return r; } };
447 
448 template <> struct vector_type<float,1> { typedef float type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const float i1) { return i1; } };
449 template <> struct vector_type<float,2> { typedef float2 type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const float i1, const float i2) { return make_float2(i1,i2); } };
450 template <> struct vector_type<float,3> { typedef float3 type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const float i1, const float i2, const float i3) { return make_float3(i1,i2,i3); } };
451 template <> struct vector_type<float,4> { typedef float4 type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const float i1, const float i2, const float i3, const float i4) { return make_float4(i1,i2,i3,i4); } };
452 
453 template <> struct vector_type<double,1> { typedef double type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const double i1) { return i1; } };
454 template <> struct vector_type<double,2> { typedef double2 type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const double i1, const double i2) { return make_double2(i1,i2); } };
455 template <> struct vector_type<double,3> { typedef double3 type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const double i1, const double i2, const double i3) { return make_double3(i1,i2,i3); } };
456 template <> struct vector_type<double,4> { typedef double4 type; CUGAR_FORCEINLINE CUGAR_HOST_DEVICE static type make(const double i1, const double i2, const double i3, const double i4) { return make_double4(i1,i2,i3,i4); } };
457 
458 template <typename T> CUGAR_FORCEINLINE CUGAR_HOST_DEVICE typename vector_type<T,1>::type make_vector(const T i1) { return vector_type<T,1>::make( i1 ); }
459 template <typename T> CUGAR_FORCEINLINE CUGAR_HOST_DEVICE typename vector_type<T,2>::type make_vector(const T i1, const T i2) { return vector_type<T,2>::make( i1, i2 ); }
460 template <typename T> CUGAR_FORCEINLINE CUGAR_HOST_DEVICE typename vector_type<T,3>::type make_vector(const T i1, const T i2, const T i3) { return vector_type<T,3>::make( i1, i2, i3 ); }
461 template <typename T> CUGAR_FORCEINLINE CUGAR_HOST_DEVICE typename vector_type<T,4>::type make_vector(const T i1, const T i2, const T i3, const T i4) { return vector_type<T,4>::make( i1, i2, i3, i4 ); }
462 
463 template <typename T> struct vector_traits {};
464 template <> struct vector_traits<char> { typedef char value_type; const static uint32 DIM = 1; };
465 template <> struct vector_traits<unsigned char> { typedef unsigned char value_type; const static uint32 DIM = 1; };
466 template <> struct vector_traits<short> { typedef short value_type; const static uint32 DIM = 1; };
467 template <> struct vector_traits<unsigned short> { typedef unsigned short value_type; const static uint32 DIM = 1; };
468 template <> struct vector_traits<int> { typedef int value_type; const static uint32 DIM = 1; };
469 template <> struct vector_traits<unsigned int> { typedef unsigned int value_type; const static uint32 DIM = 1; };
470 template <> struct vector_traits<int64> { typedef int64 value_type; const static uint32 DIM = 1; };
471 template <> struct vector_traits<uint64> { typedef uint64 value_type; const static uint32 DIM = 1; };
472 template <> struct vector_traits<float> { typedef float value_type; const static uint32 DIM = 1; };
473 template <> struct vector_traits<double> { typedef double value_type; const static uint32 DIM = 1; };
474 template <> struct vector_traits<char2> { typedef char value_type; const static uint32 DIM = 2; };
475 template <> struct vector_traits<char3> { typedef char value_type; const static uint32 DIM = 3; };
476 template <> struct vector_traits<char4> { typedef char value_type; const static uint32 DIM = 4; };
477 template <> struct vector_traits<uchar2> { typedef unsigned char value_type; const static uint32 DIM = 2; };
478 template <> struct vector_traits<uchar3> { typedef unsigned char value_type; const static uint32 DIM = 3; };
479 template <> struct vector_traits<uchar4> { typedef unsigned char value_type; const static uint32 DIM = 4; };
480 template <> struct vector_traits<short2> { typedef short value_type; const static uint32 DIM = 2; };
481 template <> struct vector_traits<short3> { typedef short value_type; const static uint32 DIM = 3; };
482 template <> struct vector_traits<short4> { typedef short value_type; const static uint32 DIM = 4; };
483 template <> struct vector_traits<ushort2> { typedef unsigned short value_type; const static uint32 DIM = 2; };
484 template <> struct vector_traits<ushort3> { typedef unsigned short value_type; const static uint32 DIM = 3; };
485 template <> struct vector_traits<ushort4> { typedef unsigned short value_type; const static uint32 DIM = 4; };
486 template <> struct vector_traits<int2> { typedef int value_type; const static uint32 DIM = 2; };
487 template <> struct vector_traits<int3> { typedef int value_type; const static uint32 DIM = 3; };
488 template <> struct vector_traits<int4> { typedef int value_type; const static uint32 DIM = 4; };
489 template <> struct vector_traits<uint2> { typedef unsigned int value_type; const static uint32 DIM = 2; };
490 template <> struct vector_traits<uint3> { typedef unsigned int value_type; const static uint32 DIM = 3; };
491 template <> struct vector_traits<uint4> { typedef unsigned int value_type; const static uint32 DIM = 4; };
492 template <> struct vector_traits<float2> { typedef float value_type; const static uint32 DIM = 2; };
493 template <> struct vector_traits<float3> { typedef float value_type; const static uint32 DIM = 3; };
494 template <> struct vector_traits<float4> { typedef float value_type; const static uint32 DIM = 4; };
495 template <> struct vector_traits<double2> { typedef float value_type; const static uint32 DIM = 2; };
496 template <> struct vector_traits<uint64_2> { typedef uint64 value_type; const static uint32 DIM = 2; };
497 template <> struct vector_traits<uint64_3> { typedef uint64 value_type; const static uint32 DIM = 3; };
498 template <> struct vector_traits<uint64_4> { typedef uint64 value_type; const static uint32 DIM = 4; };
499 template <> struct vector_traits<int64_2> { typedef int64 value_type; const static uint32 DIM = 2; };
500 template <> struct vector_traits<int64_3> { typedef int64 value_type; const static uint32 DIM = 3; };
501 template <> struct vector_traits<int64_4> { typedef int64 value_type; const static uint32 DIM = 4; };
502 
504 
507 template <typename T>
508 inline CUGAR_HOST_DEVICE T sgn(const T x) { return x > 0 ? T(1) : T(-1); }
509 
512 inline CUGAR_HOST_DEVICE float round(const float x)
513 {
514  const int y = x > 0.0f ? int(x) : int(x)-1;
515  return (x - float(y) > 0.5f) ? float(y)+1.0f : float(y);
516 }
517 
520 inline CUGAR_HOST_DEVICE int32 abs(const int32 a) { return a < 0 ? -a : a; }
521 
524 inline CUGAR_HOST_DEVICE int64 abs(const int64 a) { return a < 0 ? -a : a; }
525 
528 inline CUGAR_HOST_DEVICE float abs(const float a) { return fabsf(a); }
529 
532 inline CUGAR_HOST_DEVICE double abs(const double a) { return fabs(a); }
533 
536 inline CUGAR_HOST_DEVICE float min(const float a, const float b) { return a < b ? a : b; }
537 
540 inline CUGAR_HOST_DEVICE float max(const float a, const float b) { return a > b ? a : b; }
541 
544 inline CUGAR_HOST_DEVICE int8 min(const int8 a, const int8 b) { return a < b ? a : b; }
545 
548 inline CUGAR_HOST_DEVICE int8 max(const int8 a, const int8 b) { return a > b ? a : b; }
549 
552 inline CUGAR_HOST_DEVICE uint8 min(const uint8 a, const uint8 b) { return a < b ? a : b; }
553 
556 inline CUGAR_HOST_DEVICE uint8 max(const uint8 a, const uint8 b) { return a > b ? a : b; }
557 
560 inline CUGAR_HOST_DEVICE uint16 min(const uint16 a, const uint16 b) { return a < b ? a : b; }
561 
564 inline CUGAR_HOST_DEVICE uint16 max(const uint16 a, const uint16 b) { return a > b ? a : b; }
565 
568 inline CUGAR_HOST_DEVICE int32 min(const int32 a, const int32 b) { return a < b ? a : b; }
569 
572 inline CUGAR_HOST_DEVICE int32 max(const int32 a, const int32 b) { return a > b ? a : b; }
573 
576 inline CUGAR_HOST_DEVICE uint32 min(const uint32 a, const uint32 b) { return a < b ? a : b; }
577 
580 inline CUGAR_HOST_DEVICE uint32 max(const uint32 a, const uint32 b) { return a > b ? a : b; }
581 
584 inline CUGAR_HOST_DEVICE int64 min(const int64 a, const int64 b) { return a < b ? a : b; }
585 
588 inline CUGAR_HOST_DEVICE int64 max(const int64 a, const int64 b) { return a > b ? a : b; }
589 
592 inline CUGAR_HOST_DEVICE uint64 min(const uint64 a, const uint64 b) { return a < b ? a : b; }
593 
596 inline CUGAR_HOST_DEVICE uint64 max(const uint64 a, const uint64 b) { return a > b ? a : b; }
597 
600 inline CUGAR_HOST_DEVICE uint32 quantize(const float x, const uint32 n)
601 {
602  return (uint32)max( min( int32( x * float(n) ), int32(n-1) ), int32(0) );
603 }
606 inline float CUGAR_HOST_DEVICE mod(const float x, const float m) { return x > 0.0f ? fmodf( x, m ) : m - fmodf( -x, m ); }
607 
610 inline float CUGAR_HOST_DEVICE sqr(const float x) { return x*x; }
611 
614 inline double CUGAR_HOST_DEVICE sqr(const double x) { return x*x; }
615 
618 inline CUGAR_HOST_DEVICE uint32 log2(uint32 n)
619 {
620  unsigned int c = 0;
621  if (n & 0xffff0000u) { n >>= 16; c |= 16; }
622  if (n & 0xff00) { n >>= 8; c |= 8; }
623  if (n & 0xf0) { n >>= 4; c |= 4; }
624  if (n & 0xc) { n >>= 2; c |= 2; }
625  if (n & 0x2) c |= 1;
626  return c;
627 /* uint32 m = 0;
628  while (n > 0)
629  {
630  n >>= 1;
631  m++;
632  }
633  return m-1;*/
634 }
635 
636 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE
637 float saturate(const float x)
638 {
639 #ifdef CUGAR_DEVICE_COMPILATION
640  return ::saturate(x);
641 #else
642  return max( min( x, 1.0f ), 0.0f );
643 #endif
644 }
645 
648 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE
649 uint32 hash(uint32 a)
650 {
651  a = (a+0x7ed55d16) + (a<<12);
652  a = (a^0xc761c23c) ^ (a>>19);
653  a = (a+0x165667b1) + (a<<5);
654  a = (a+0xd3a2646c) ^ (a<<9);
655  a = (a+0xfd7046c5) + (a<<3);
656  a = (a^0xb55a4f09) ^ (a>>16);
657  return a;
658 }
659 
662 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE
663 uint32 hash2(uint32 key)
664 {
665  key += ~(key << 15);
666  key ^= (key >> 10);
667  key += (key << 3);
668  key ^= (key >> 6);
669  key += ~(key << 11);
670  key ^= (key >> 16);
671  return key;
672 }
673 
676 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE
677 uint64 hash(uint64 key)
678 {
679  key += ~(key << 32);
680  key ^= (key >> 22);
681  key += ~(key << 13);
682  key ^= (key >> 8);
683  key += (key << 3);
684  key ^= (key >> 15);
685  key += ~(key << 27);
686  key ^= (key >> 31);
687  return key;
688 }
689 
692 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE
693 uint64 hash2(uint64 key)
694 {
695  return (key >> 32) ^ key;
696 }
697 
700 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE
701 uint64 hash3(uint64 key)
702 {
703  uint32 hash = 0u;
704 
705  #if defined(__CUDA_ARCH__)
706  #pragma unroll
707  #endif
708  for (uint32 i = 0; i < 8; ++i)
709  {
710  hash = (hash << 4) + ((key >> (i*8)) & 255u); // shift/mix
711 
712  // get high nybble
713  const uint32 hi_bits = hash & 0xF0000000;
714  if (hi_bits != 0u)
715  hash ^= hi_bits >> 24; // xor high nybble with second nybble
716 
717  hash &= ~hi_bits; // clear high nybble
718  }
719  return hash;
720 }
721 
722 #define CUGAR_RAND_A 1664525
723 #define CUGAR_RAND_C 1013904223
724 
727 struct LCG_random
728 {
729  static const uint32 MAX = 0xFFFFFFFF;
730 
731  CUGAR_FORCEINLINE CUGAR_HOST_DEVICE LCG_random(const uint32 s = 0) : m_s(s) {}
732 
733  CUGAR_FORCEINLINE CUGAR_HOST_DEVICE uint32 next() { m_s = m_s*CUGAR_RAND_A + CUGAR_RAND_C; return m_s; }
734 
735  uint32 m_s;
736 };
737 
740 struct FLCG_random : LCG_random
741 {
742  CUGAR_FORCEINLINE CUGAR_HOST_DEVICE FLCG_random(const uint32 s = 0) : LCG_random(s) {}
743 
744  CUGAR_FORCEINLINE CUGAR_HOST_DEVICE float next() { return this->LCG_random::next() / float(LCG_random::MAX); }
745 };
746 
752 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE
753 float randfloat(unsigned i, unsigned p)
754 {
755  i ^= p;
756  i ^= i >> 17;
757  i ^= i >> 10; i *= 0xb36534e5;
758  i ^= i >> 12;
759  i ^= i >> 21; i *= 0x93fc4795;
760  i ^= 0xdf6e307f;
761  i ^= i >> 17; i *= 1 | p >> 18;
762  return i * (1.0f / 4294967808.0f);
763 }
764 
767 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE
768 uint32 reverse_bits(const uint32 n)
769 {
770  uint32 bits = n;
771  bits = (bits << 16u) | (bits >> 16u);
772  bits = ((bits & 0x55555555u) << 1u) | ((bits & 0xAAAAAAAAu) >> 1u);
773  bits = ((bits & 0x33333333u) << 2u) | ((bits & 0xCCCCCCCCu) >> 2u);
774  bits = ((bits & 0x0F0F0F0Fu) << 4u) | ((bits & 0xF0F0F0F0u) >> 4u);
775  bits = ((bits & 0x00FF00FFu) << 8u) | ((bits & 0xFF00FF00u) >> 8u);
776  return bits;
777 }
778 
781 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE
782 float radical_inverse(unsigned int n)
783 {
784 #if 0
785  double result = 0.0;
786  unsigned int remainder;
787  unsigned int m, bj = 1;
788 
789  const unsigned int b = 2u;
790 
791  do
792  {
793  bj *= b;
794  m = n;
795  n /= b;
796 
797  remainder = m - n * b;
798 
799  result += double(remainder) / double(bj);
800  } while (n > 0);
801 
802  return float(result);
803 #else
804  uint32 bits = n;
805  bits = (bits << 16u) | (bits >> 16u);
806  bits = ((bits & 0x55555555u) << 1u) | ((bits & 0xAAAAAAAAu) >> 1u);
807  bits = ((bits & 0x33333333u) << 2u) | ((bits & 0xCCCCCCCCu) >> 2u);
808  bits = ((bits & 0x0F0F0F0Fu) << 4u) | ((bits & 0xF0F0F0F0u) >> 4u);
809  bits = ((bits & 0x00FF00FFu) << 8u) | ((bits & 0xFF00FF00u) >> 8u);
810  return float(bits) * 2.3283064365386963e-10f; // / 0x100000000
811 #endif
812 };
813 
816 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE
817 uint64 radical_inverse(uint64 bits, const uint64 scramble)
818 {
819  bits = (bits << 32) | (bits >> 32);
820  bits = ((bits & 0x0000ffff0000ffffULL) << 16) |
821  ((bits & 0xffff0000ffff0000ULL) >> 16);
822  bits = ((bits & 0x00ff00ff00ff00ffULL) << 8) |
823  ((bits & 0xff00ff00ff00ff00ULL) >> 8);
824  bits = ((bits & 0x0f0f0f0f0f0f0f0fULL) << 4) |
825  ((bits & 0xf0f0f0f0f0f0f0f0ULL) >> 4);
826  bits = ((bits & 0x3333333333333333ULL) << 2) |
827  ((bits & 0xccccccccccccccccULL) >> 2);
828  bits = ((bits & 0x5555555555555555ULL) << 1) |
829  ((bits & 0xaaaaaaaaaaaaaaaaULL) >> 1);
830  return (scramble ^ bits) >> (64 - 52); // Account for 52 bits precision.
831 }
832 
835 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE
836 uint32 permute(uint32 i, uint32 l, uint32 p)
837 {
838  uint32 w = l - 1;
839  w |= w >> 1;
840  w |= w >> 2;
841  w |= w >> 4;
842  w |= w >> 8;
843  w |= w >> 16;
844  do {
845  i ^= p; i *= 0xe170893d;
846  i ^= p >> 16;
847  i ^= (i & w) >> 4;
848  i ^= p >> 8; i *= 0x0929eb3f;
849  i ^= p >> 23;
850  i ^= (i & w) >> 1; i *= 1 | p >> 27;
851  i *= 0x6935fa69;
852  i ^= (i & w) >> 11; i *= 0x74dcb303;
853  i ^= (i & w) >> 2; i *= 0x9e501cc3;
854  i ^= (i & w) >> 2; i *= 0xc860a3df;
855  i &= w;
856  i ^= i >> 5;
857  } while (i >= l);
858 
859  return (i + p) % l;
860 }
861 
862 #if defined(__CUDA_ARCH__)
863 
864 CUGAR_FORCEINLINE CUGAR_DEVICE
865 uint8 min3(const uint8 op1, const uint8 op2, const uint8 op3)
866 {
867  uint32 r;
868  asm( " vmin.u32.u32.u32.min %0, %1, %2, %3;" : "=r"(r) : "r"(uint32(op1)), "r"(uint32(op2)), "r"(uint32(op3)) );
869  return r;
870 }
871 
872 CUGAR_FORCEINLINE CUGAR_DEVICE
873 uint32 min3(const uint32 op1, const uint32 op2, const uint32 op3)
874 {
875  uint32 r;
876  asm( " vmin.u32.u32.u32.min %0, %1, %2, %3;" : "=r"(r) : "r"(op1), "r"(op2), "r"(op3) );
877  return r;
878 }
879 
880 CUGAR_FORCEINLINE CUGAR_DEVICE
881 uint32 max3(const uint32 op1, const uint32 op2, const uint32 op3)
882 {
883  uint32 r;
884  asm( " vmax.u32.u32.u32.max %0, %1, %2, %3;" : "=r"(r) : "r"(op1), "r"(op2), "r"(op3) );
885  return r;
886 }
887 
888 CUGAR_FORCEINLINE CUGAR_DEVICE
889 int32 min3(const int32 op1, const int32 op2, const int32 op3)
890 {
891  uint32 r;
892  asm( " vmin.s32.s32.s32.min %0, %1, %2, %3;" : "=r"(r) : "r"(op1), "r"(op2), "r"(op3) );
893  return r;
894 }
895 
896 CUGAR_FORCEINLINE CUGAR_DEVICE
897 int32 max3(const int32 op1, const int32 op2, const int32 op3)
898 {
899  uint32 r;
900  asm( " vmax.s32.s32.s32.max %0, %1, %2, %3;" : "=r"(r) : "r"(op1), "r"(op2), "r"(op3) );
901  return r;
902 }
903 
904 #else
905 
906 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE
907 uint8 min3(const uint8 op1, const uint8 op2, const uint8 op3)
908 {
909  return cugar::min( op1, cugar::min( op2, op3 ) );
910 }
911 
912 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE
913 uint32 min3(const uint32 op1, const uint32 op2, const uint32 op3)
914 {
915  return cugar::min( op1, cugar::min( op2, op3 ) );
916 }
917 
918 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE
919 uint32 max3(const uint32 op1, const uint32 op2, const uint32 op3)
920 {
921  return cugar::max( op1, cugar::max( op2, op3 ) );
922 }
923 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE
924 int32 min3(const int32 op1, const int32 op2, const int32 op3)
925 {
926  return cugar::min( op1, cugar::min( op2, op3 ) );
927 }
928 
929 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE
930 int32 max3(const int32 op1, const int32 op2, const int32 op3)
931 {
932  return cugar::max( op1, cugar::max( op2, op3 ) );
933 }
934 
935 #endif
936 
937 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE
938 float min3(const float op1, const float op2, const float op3)
939 {
940  return cugar::min( op1, cugar::min( op2, op3 ) );
941 }
942 
943 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE
944 float max3(const float op1, const float op2, const float op3)
945 {
946  return cugar::max( op1, cugar::max( op2, op3 ) );
947 }
948 
949 #ifdef __CUDA_ARCH__
950 
951 inline CUGAR_DEVICE float fast_pow(const float a, const float b)
952 {
953  return __powf(a,b);
954 }
955 inline CUGAR_DEVICE float fast_sin(const float x)
956 {
957  return __sinf(x);
958 }
959 inline CUGAR_DEVICE float fast_cos(const float x)
960 {
961  return __cosf(x);
962 }
963 inline CUGAR_DEVICE float fast_sqrt(const float x)
964 {
965  return __fsqrt_rn(x);
966 }
967 
968 #else
969 
970 inline CUGAR_HOST_DEVICE float fast_pow(const float a, const float b)
971 {
972  return ::powf(a,b);
973 }
974 inline CUGAR_HOST_DEVICE float fast_sin(const float x)
975 {
976  return sinf(x);
977 }
978 inline CUGAR_HOST_DEVICE float fast_cos(const float x)
979 {
980  return cosf(x);
981 }
982 inline CUGAR_HOST_DEVICE float fast_sqrt(const float x)
983 {
984  return sqrtf(x);
985 }
986 
987 #endif
988 
989 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE
990 void sincosf(float phi, float* s, float* c)
991 {
992 #if defined(CUGAR_DEVIE_COMPILATION)
993  ::sincosf(phi, s, c);
994 #else
995  *s = sinf(phi);
996  *c = cosf(phi);
997 #endif
998 }
999 
1000 CUGAR_FORCEINLINE CUGAR_HOST_DEVICE
1001 float rsqrtf(float x)
1002 {
1003 #if defined(CUGAR_DEVIE_COMPILATION)
1004  return ::rsqrtf(x);
1005 #else
1006  return 1.0f / sqrtf(x);
1007 #endif
1008 }
1009 
1010 #ifdef __CUDACC__
1011 inline CUGAR_DEVICE uint16 float_to_half(const float x) { return __float2half_rn(x); }
1012 inline CUGAR_DEVICE float half_to_float(const uint32 h) { return __half2float(h); }
1013 #endif
1014 
1019 template <typename T>
1020 struct field_traits
1021 {
1022 #ifdef __CUDACC__
1023  CUGAR_HOST_DEVICE static T min() { return T(); }
1026 
1029  CUGAR_HOST_DEVICE static T max() { return T(); }
1030 #else
1031  static T min()
1034  {
1035  return std::numeric_limits<T>::is_integer ?
1036  std::numeric_limits<T>::min() :
1037  -std::numeric_limits<T>::max();
1038  }
1041  static T max() { return std::numeric_limits<T>::max(); }
1042 #endif
1043 };
1044 
1047 template <>
1048 struct field_traits<int8>
1049 {
1050  CUGAR_HOST_DEVICE static int8 min() { return -128; }
1051  CUGAR_HOST_DEVICE static int8 max() { return 127; }
1052 };
1055 template <>
1056 struct field_traits<int16>
1057 {
1058  CUGAR_HOST_DEVICE static int16 min() { return -32768; }
1059  CUGAR_HOST_DEVICE static int16 max() { return 32767; }
1060 };
1063 template <>
1064 struct field_traits<int32>
1065 {
1066  CUGAR_HOST_DEVICE static int32 min() { return -(1 << 30); }
1067  CUGAR_HOST_DEVICE static int32 max() { return (1 << 30); }
1068 };
1071 template <>
1072 struct field_traits<int64>
1073 {
1074  CUGAR_HOST_DEVICE static int64 min() { return -(int64(1) << 62); }
1075  CUGAR_HOST_DEVICE static int64 max() { return (int64(1) << 62); }
1076 };
1077 
1078 #ifdef __CUDACC__
1079 template <>
1082 struct field_traits<float>
1083 {
1084  CUGAR_HOST_DEVICE static float min() { return -float(1.0e+30f); }
1085  CUGAR_HOST_DEVICE static float max() { return float(1.0e+30f); }
1086 };
1089 template <>
1090 struct field_traits<double>
1091 {
1092  CUGAR_HOST_DEVICE static double min() { return -double(1.0e+30); }
1093  CUGAR_HOST_DEVICE static double max() { return double(1.0e+30); }
1094 };
1097 template <>
1098 struct field_traits<uint32>
1099 {
1100  CUGAR_HOST_DEVICE static uint32 min() { return 0; }
1101  CUGAR_HOST_DEVICE static uint32 max() { return uint32(-1); }
1102 };
1105 template <>
1106 struct field_traits<uint64>
1107 {
1108  CUGAR_HOST_DEVICE static uint64 min() { return 0; }
1109  CUGAR_HOST_DEVICE static uint64 max() { return uint64(-1); }
1110 };
1111 #endif
1112 
1113 using ::sinf;
1114 using ::sin;
1115 using ::cosf;
1116 using ::cos;
1117 using ::sqrtf;
1118 using ::sqrt;
1119 using ::expf;
1120 using ::exp;
1121 using ::logf;
1122 using ::log;
1123 
1126 
1127 } // namespace cugar
cugar::round_z
CUGAR_HOST_DEVICE L round_z(const L x, const R y)
Definition: numbers.h:204
cugar::quantize
CUGAR_HOST_DEVICE uint32 quantize(const float x, const uint32 n)
Definition: numbers.h:600
cugar::radical_inverse
CUGAR_FORCEINLINE CUGAR_HOST_DEVICE float radical_inverse(unsigned int n)
Definition: numbers.h:782
cugar::vector4_storage
Definition: numbers.h:395
cugar::mod
float CUGAR_HOST_DEVICE mod(const float x, const float m)
Definition: numbers.h:606
cugar::begin
thrust::device_vector< T >::iterator begin(thrust::device_vector< T > &vec)
Definition: thrust_view.h:89
cugar::vector1_storage
Definition: numbers.h:371
cugar::LCG_random
Definition: numbers.h:727
cugar::vector3_storage
Definition: numbers.h:387
cugar::log2
CUGAR_HOST_DEVICE uint32 log2(uint32 n)
Definition: numbers.h:618
cugar::min
CUGAR_HOST_DEVICE uint64 min(const uint64 a, const uint64 b)
Definition: numbers.h:592
cugar::vector2_storage
Definition: numbers.h:379
cugar::permute
CUGAR_FORCEINLINE CUGAR_HOST_DEVICE uint32 permute(uint32 i, uint32 l, uint32 p)
Definition: numbers.h:836
cugar::round_i
CUGAR_HOST_DEVICE L round_i(const L x, const R y)
Definition: numbers.h:198
cugar::sgn
CUGAR_HOST_DEVICE T sgn(const T x)
Definition: numbers.h:508
cugar::vector_traits
Definition: numbers.h:463
cugar::divide_ri
CUGAR_HOST_DEVICE L divide_ri(const L x, const R y)
Definition: numbers.h:180
cugar::FLCG_random
Definition: numbers.h:740
cugar::hash2
CUGAR_FORCEINLINE CUGAR_HOST_DEVICE uint32 hash2(uint32 key)
Definition: numbers.h:663
cugar::sqr
float CUGAR_HOST_DEVICE sqr(const float x)
Definition: numbers.h:610
cugar::field_traits
Definition: numbers.h:1020
cugar::max
CUGAR_HOST_DEVICE uint64 max(const uint64 a, const uint64 b)
Definition: numbers.h:596
cugar::abs
CUGAR_HOST_DEVICE int32 abs(const int32 a)
Definition: numbers.h:520
cugar::count_occurrences
CUGAR_FORCEINLINE CUGAR_HOST_DEVICE uint32 count_occurrences(const Iterator begin, uint32 size, const T val, const uint32 max_occ=uint32(-1))
Definition: numbers.h:162
cugar::vector_type
Definition: numbers.h:368
cugar::lo_bits
CUGAR_FORCEINLINE CUGAR_HOST_DEVICE uint32 lo_bits()
Definition: numbers.h:148
cugar::hi_bits
CUGAR_FORCEINLINE CUGAR_HOST_DEVICE uint32 hi_bits()
Definition: numbers.h:155
cugar
Define a vector_view POD type and plain_view() for std::vector.
Definition: diff.h:38
cugar::field_traits::max
static T max()
Definition: numbers.h:1041
cugar::divide_rz
CUGAR_HOST_DEVICE L divide_rz(const L x, const R y)
Definition: numbers.h:189
cugar::binary_cast
CUGAR_FORCEINLINE CUGAR_HOST_DEVICE Out binary_cast(const In in)
Definition: types.h:288
cugar::comp
CUGAR_FORCEINLINE CUGAR_HOST_DEVICE uint8 comp(const uchar2 a, const char c)
Definition: numbers.h:218
cugar::reverse_bits
CUGAR_FORCEINLINE CUGAR_HOST_DEVICE uint32 reverse_bits(const uint32 n)
Definition: numbers.h:768
cugar::hash
CUGAR_FORCEINLINE CUGAR_HOST_DEVICE uint32 hash(uint32 a)
Definition: numbers.h:649
cugar::round
CUGAR_HOST_DEVICE L round(const L x, const R y)
Definition: numbers.h:210
cugar::randfloat
CUGAR_FORCEINLINE CUGAR_HOST_DEVICE float randfloat(unsigned i, unsigned p)
Definition: numbers.h:753
cugar::hash3
CUGAR_FORCEINLINE CUGAR_HOST_DEVICE uint64 hash3(uint64 key)
Definition: numbers.h:701
Generated by   doxygen 1.8.13