Ctrl+K

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-rocrand/checkouts/latest/library/include/rocrand/rocrand_sobol64.h Source File

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-rocrand/checkouts/latest/library/include/rocrand/rocrand_sobol64.h Source File#

API library: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-rocrand/checkouts/latest/library/include/rocrand/rocrand_sobol64.h Source File
API library
rocrand_sobol64.h
1 // Copyright (c) 2021-2022 Advanced Micro Devices, Inc. All rights reserved.
2 //
3 // Permission is hereby granted, free of charge, to any person obtaining a copy
4 // of this software and associated documentation files (the "Software"), to deal
5 // in the Software without restriction, including without limitation the rights
6 // to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
7 // copies of the Software, and to permit persons to whom the Software is
8 // furnished to do so, subject to the following conditions:
9 //
10 // The above copyright notice and this permission notice shall be included in
11 // all copies or substantial portions of the Software.
12 //
13 // THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
14 // IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
15 // FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
16 // AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
17 // LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
18 // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
19 // THE SOFTWARE.
20 
21 #ifndef ROCRAND_SOBOL64_H_
22 #define ROCRAND_SOBOL64_H_
23 
24 #ifndef FQUALIFIERS
25 #define FQUALIFIERS __forceinline__ __device__
26 #endif // FQUALIFIERS_
27 
28 #include "rocrand/rocrand_common.h"
29 
30 namespace rocrand_device {
31 
32 template<bool UseSharedVectors>
33 struct sobol64_state
34 {
35  unsigned long long int d;
36  unsigned long long int i;
37  unsigned long long int vectors[64];
38 
39  FQUALIFIERS
40  sobol64_state() : d(), i(), vectors() { }
41 
42  FQUALIFIERS
43  sobol64_state(const unsigned long long int d,
44  const unsigned long long int i,
45  const unsigned long long int * vectors)
46  : d(d), i(i)
47  {
48  for(int k = 0; k < 64; k++)
49  {
50  this->vectors[k] = vectors[k];
51  }
52  }
53 };
54 
55 template<>
56 struct sobol64_state<true>
57 {
58  unsigned long long int d;
59  unsigned long long int i;
60  const unsigned long long int * vectors;
61 
62  FQUALIFIERS
63  sobol64_state() : d(), i(), vectors() { }
64 
65  FQUALIFIERS
66  sobol64_state(const unsigned long long int d,
67  const unsigned long long int i,
68  const unsigned long long int * vectors)
69  : d(d), i(i), vectors(vectors) { }
70 };
71 
72 template<bool UseSharedVectors>
73 class sobol64_engine
74 {
75 public:
76 
77  typedef struct sobol64_state<UseSharedVectors> sobol64_state;
78 
79  FQUALIFIERS
80  sobol64_engine() { }
81 
82  FQUALIFIERS
83  sobol64_engine(const unsigned long long int* vectors, const unsigned long long int offset)
84  : m_state(0, 0, vectors)
85  {
86  discard_state(offset);
87  }
88 
90  FQUALIFIERS
91  void discard(unsigned long long int offset)
92  {
93  discard_state(offset);
94  }
95 
96  FQUALIFIERS
97  void discard()
98  {
99  discard_state();
100  }
101 
103  FQUALIFIERS
104  void discard_stride(unsigned long long int stride)
105  {
106  discard_state_power2(stride);
107  }
108 
109  FQUALIFIERS
110  unsigned long long int operator()()
111  {
112  return this->next();
113  }
114 
115  FQUALIFIERS
116  unsigned long long int next()
117  {
118  unsigned long long int p = m_state.d;
119  discard_state();
120  return p;
121  }
122 
123  FQUALIFIERS
124  unsigned long long int current() const
125  {
126  return m_state.d;
127  }
128 
129 protected:
130  // Advances the internal state by offset times.
131  FQUALIFIERS
132  void discard_state(unsigned long long int offset)
133  {
134  m_state.i += offset;
135  const unsigned long long int g = m_state.i ^ (m_state.i >> 1ull);
136  m_state.d = 0;
137  for(int i = 0; i < 64; i++)
138  {
139  m_state.d ^= (g & (1ull << i) ? m_state.vectors[i] : 0ull);
140  }
141  }
142 
143  // Advances the internal state to the next state
144  FQUALIFIERS
145  void discard_state()
146  {
147  m_state.d ^= m_state.vectors[rightmost_zero_bit(m_state.i)];
148  m_state.i++;
149  }
150 
151  FQUALIFIERS
152  void discard_state_power2(unsigned long long int stride)
153  {
154  // Leap frog
155  //
156  // T Bradley, J Toit, M Giles, R Tong, P Woodhams
157  // Parallelisation Techniques for Random Number Generators
158  // GPU Computing Gems, 2011
159  //
160  // For power of 2 jumps only 2 bits in Gray code change values
161  // All bits lower than log2(stride) flip 2, 4... times, i.e.
162  // do not change their values.
163 
164  // log2(stride) bit
165  m_state.d ^= m_state.vectors[rightmost_zero_bit(~stride) - 1];
166  // the rightmost zero bit of i, not including the lower log2(stride) bits
167  m_state.d ^= m_state.vectors[rightmost_zero_bit(m_state.i | (stride - 1))];
168  m_state.i += stride;
169  }
170 
171  // Returns the index of the rightmost zero bit in the binary expansion of
172  // x (Gray code of the current element's index)
173  // NOTE changing unsigned long long int to unit64_t will cause compile failure on device
174  FQUALIFIERS
175  unsigned int rightmost_zero_bit(unsigned long long int x)
176  {
177  #if defined(__HIP_DEVICE_COMPILE__)
178  unsigned int z = __ffsll(~x);
179  return z ? z - 1 : 0;
180  #else
181  if(x == 0)
182  return 0;
183  unsigned long long int y = x;
184  unsigned long long int z = 1;
185  while(y & 1)
186  {
187  y >>= 1;
188  z++;
189  }
190  return z - 1;
191  #endif
192  }
193 
194 protected:
195  // State
196  sobol64_state m_state;
197 
198 }; // sobol64_engine class
199 
200 } // end namespace rocrand_device
201 
208 typedef rocrand_device::sobol64_engine<false> rocrand_state_sobol64;
210 
221 FQUALIFIERS
222 void rocrand_init(const unsigned long long int * vectors,
223  const unsigned int offset,
224  rocrand_state_sobol64 * state)
225 {
226  *state = rocrand_state_sobol64(vectors, offset);
227 }
228 
241 FQUALIFIERS
242 unsigned long long int rocrand(rocrand_state_sobol64 * state)
243 {
244  return state->next();
245 }
246 
255 FQUALIFIERS
256 void skipahead(unsigned long long int offset, rocrand_state_sobol64* state)
257 {
258  return state->discard(offset);
259 }
260  // end of group rocranddevice
262 
263 #endif // ROCRAND_sobol64_H_
rocrand
FQUALIFIERS unsigned long long int rocrand(rocrand_state_sobol64 *state)
Returns uniformly distributed random unsigned int value from [0; 2^64 - 1] range.
Definition: rocrand_sobol64.h:242
rocrand_init
FQUALIFIERS void rocrand_init(const unsigned long long int *vectors, const unsigned int offset, rocrand_state_sobol64 *state)
Initialize sobol64 state.
Definition: rocrand_sobol64.h:222
skipahead
FQUALIFIERS void skipahead(unsigned long long int offset, rocrand_state_sobol64 *state)
Updates sobol64 state to skip ahead by offset elements.
Definition: rocrand_sobol64.h:256
FQUALIFIERS
#define FQUALIFIERS
Shorthand for commonly used function qualifiers.
Definition: rocrand_uniform.h:31
Generated by doxygen 1.9.1