1 /*
2 * Copyright (c) 2016, NVIDIA CORPORATION. All rights reserved.
3 *
4 * Redistribution and use in source and binary forms, with or without
5 * modification, are permitted provided that the following conditions
6 * are met:
7 * * Redistributions of source code must retain the above copyright
8 * notice, this list of conditions and the following disclaimer.
9 * * Redistributions in binary form must reproduce the above copyright
10 * notice, this list of conditions and the following disclaimer in the
11 * documentation and/or other materials provided with the distribution.
12 * * Neither the name of NVIDIA CORPORATION nor the names of its
13 * contributors may be used to endorse or promote products derived
14 * from this software without specific prior written permission.
15 *
16 * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
17 * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
18 * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
19 * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR
20 * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
21 * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
22 * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
23 * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
24 * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
25 * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
26 * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
27 */
28
29 #include <optixu/optixu_math_namespace.h>
30
31 template<unsigned int N>
tea(unsigned int val0,unsigned int val1)32 static __host__ __device__ __inline__ unsigned int tea( unsigned int val0, unsigned int val1 )
33 {
34 unsigned int v0 = val0;
35 unsigned int v1 = val1;
36 unsigned int s0 = 0;
37
38 for( unsigned int n = 0; n < N; n++ )
39 {
40 s0 += 0x9e3779b9;
41 v0 += ((v1<<4)+0xa341316c)^(v1+s0)^((v1>>5)+0xc8013ea4);
42 v1 += ((v0<<4)+0xad90777d)^(v0+s0)^((v0>>5)+0x7e95761e);
43 }
44
45 return v0;
46 }
47
48 // Generate random unsigned int in [0, 2^24)
lcg(unsigned int & prev)49 static __host__ __device__ __inline__ unsigned int lcg(unsigned int &prev)
50 {
51 const unsigned int LCG_A = 1664525u;
52 const unsigned int LCG_C = 1013904223u;
53 prev = (LCG_A * prev + LCG_C);
54 return prev & 0x00FFFFFF;
55 }
56
lcg2(unsigned int & prev)57 static __host__ __device__ __inline__ unsigned int lcg2(unsigned int &prev)
58 {
59 prev = (prev*8121 + 28411) % 134456;
60 return prev;
61 }
62
63 // Generate random float in [0, 1)
rnd(unsigned int & prev)64 static __host__ __device__ __inline__ float rnd(unsigned int &prev)
65 {
66 return ((float) lcg(prev) / (float) 0x01000000);
67 }
68
69 // Multiply with carry
mwc()70 static __host__ __inline__ unsigned int mwc()
71 {
72 static unsigned long long r[4];
73 static unsigned long long carry;
74 static bool init = false;
75 if( !init ) {
76 init = true;
77 unsigned int seed = 7654321u, seed0, seed1, seed2, seed3;
78 r[0] = seed0 = lcg2(seed);
79 r[1] = seed1 = lcg2(seed0);
80 r[2] = seed2 = lcg2(seed1);
81 r[3] = seed3 = lcg2(seed2);
82 carry = lcg2(seed3);
83 }
84
85 unsigned long long sum = 2111111111ull * r[3] +
86 1492ull * r[2] +
87 1776ull * r[1] +
88 5115ull * r[0] +
89 1ull * carry;
90 r[3] = r[2];
91 r[2] = r[1];
92 r[1] = r[0];
93 r[0] = static_cast<unsigned int>(sum); // lower half
94 carry = static_cast<unsigned int>(sum >> 32); // upper half
95 return static_cast<unsigned int>(r[0]);
96 }
97
random1u()98 static __host__ __inline__ unsigned int random1u()
99 {
100 #if 0
101 return rand();
102 #else
103 return mwc();
104 #endif
105 }
106
random2u()107 static __host__ __inline__ optix::uint2 random2u()
108 {
109 return optix::make_uint2(random1u(), random1u());
110 }
111
fillRandBuffer(unsigned int * seeds,unsigned int N)112 static __host__ __inline__ void fillRandBuffer( unsigned int *seeds, unsigned int N )
113 {
114 for( unsigned int i=0; i<N; ++i )
115 seeds[i] = mwc();
116 }
117
rot_seed(unsigned int seed,unsigned int frame)118 static __host__ __device__ __inline__ unsigned int rot_seed( unsigned int seed, unsigned int frame )
119 {
120 return seed ^ frame;
121 }
122