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

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

API library: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-rocrand/checkouts/docs-6.2.1/library/include/rocrand/rocrand_sobol32.h Source File
API library
rocrand_sobol32.h
1 // Copyright (c) 2017-2024 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_SOBOL32_H_
22 #define ROCRAND_SOBOL32_H_
23 
24 #include "rocrand/rocrand_common.h"
25 
26 namespace rocrand_device {
27 
28 template<bool UseSharedVectors>
29 struct sobol32_state
30 {
31  unsigned int d;
32  unsigned int i;
33  unsigned int vectors[32];
34 
35  __forceinline__ __device__ __host__ sobol32_state() : d(), i(), vectors() {}
36 
37  __forceinline__ __device__ __host__ sobol32_state(const unsigned int d,
38  const unsigned int i,
39  const unsigned int* vectors)
40  : d(d), i(i)
41  {
42  for(int k = 0; k < 32; k++)
43  {
44  this->vectors[k] = vectors[k];
45  }
46  }
47 };
48 
49 template<>
50 struct sobol32_state<true>
51 {
52  unsigned int d;
53  unsigned int i;
54  const unsigned int * vectors;
55 
56  __forceinline__ __device__ __host__ sobol32_state() : d(), i(), vectors() {}
57 
58  __forceinline__ __device__ __host__ sobol32_state(const unsigned int d,
59  const unsigned int i,
60  const unsigned int* vectors)
61  : d(d), i(i), vectors(vectors)
62  {}
63 };
64 
65 template<bool UseSharedVectors>
66 class sobol32_engine
67 {
68 public:
69 
70  typedef struct sobol32_state<UseSharedVectors> sobol32_state;
71 
72  __forceinline__ __device__ __host__ sobol32_engine() {}
73 
74  __forceinline__ __device__ __host__ sobol32_engine(const unsigned int* vectors,
75  const unsigned int offset)
76  : m_state(0, 0, vectors)
77  {
78  discard_state(offset);
79  }
80 
82  __forceinline__ __device__ __host__ void discard(unsigned int offset)
83  {
84  discard_state(offset);
85  }
86 
87  __forceinline__ __device__ __host__ void discard()
88  {
89  discard_state();
90  }
91 
93  __forceinline__ __device__ __host__ void discard_stride(unsigned int stride)
94  {
95  discard_state_power2(stride);
96  }
97 
98  __forceinline__ __device__ __host__ unsigned int operator()()
99  {
100  return this->next();
101  }
102 
103  __forceinline__ __device__ __host__ unsigned int next()
104  {
105  unsigned int p = m_state.d;
106  discard_state();
107  return p;
108  }
109 
110  __forceinline__ __device__ __host__ unsigned int current() const
111  {
112  return m_state.d;
113  }
114 
115  __forceinline__ __device__ __host__ static constexpr bool uses_shared_vectors()
116  {
117  return UseSharedVectors;
118  }
119 
120 protected:
121  // Advances the internal state by offset times.
122  __forceinline__ __device__ __host__ void discard_state(unsigned int offset)
123  {
124  m_state.i += offset;
125  const unsigned int g = m_state.i ^ (m_state.i >> 1);
126  m_state.d = 0;
127  for(int i = 0; i < 32; i++)
128  {
129  m_state.d ^= (g & (1U << i) ? m_state.vectors[i] : 0);
130  }
131  }
132 
133  // Advances the internal state to the next state
134  __forceinline__ __device__ __host__ void discard_state()
135  {
136  m_state.d ^= m_state.vectors[rightmost_zero_bit(m_state.i)];
137  m_state.i++;
138  }
139 
140  __forceinline__ __device__ __host__ void discard_state_power2(unsigned int stride)
141  {
142  // Leap frog
143  //
144  // T Bradley, J Toit, M Giles, R Tong, P Woodhams
145  // Parallelisation Techniques for Random Number Generators
146  // GPU Computing Gems, 2011
147  //
148  // For power of 2 jumps only 2 bits in Gray code change values
149  // All bits lower than log2(stride) flip 2, 4... times, i.e.
150  // do not change their values.
151 
152  // log2(stride) bit
153  m_state.d ^= m_state.vectors[rightmost_zero_bit(~stride) - 1];
154  // the rightmost zero bit of i, not including the lower log2(stride) bits
155  m_state.d ^= m_state.vectors[rightmost_zero_bit(m_state.i | (stride - 1))];
156  m_state.i += stride;
157  }
158 
159  // Returns the index of the rightmost zero bit in the binary expansion of
160  // x (Gray code of the current element's index)
161  __forceinline__ __device__ __host__ unsigned int rightmost_zero_bit(unsigned int x)
162  {
163  #if defined(__HIP_DEVICE_COMPILE__)
164  unsigned int z = __ffs(~x);
165  return z ? z - 1 : 0;
166  #else
167  if(x == 0)
168  return 0;
169  unsigned int y = x;
170  unsigned int z = 1;
171  while(y & 1)
172  {
173  y >>= 1;
174  z++;
175  }
176  return z - 1;
177  #endif
178  }
179 
180 protected:
181  // State
182  sobol32_state m_state;
183 
184 }; // sobol32_engine class
185 
186 } // end namespace rocrand_device
187 
194 typedef rocrand_device::sobol32_engine<false> rocrand_state_sobol32;
196 
207 __forceinline__ __device__ __host__ void rocrand_init(const unsigned int* vectors,
208  const unsigned int offset,
209  rocrand_state_sobol32* state)
210 {
211  *state = rocrand_state_sobol32(vectors, offset);
212 }
213 
226 __forceinline__ __device__ __host__ unsigned int rocrand(rocrand_state_sobol32* state)
227 {
228  return state->next();
229 }
230 
239 __forceinline__ __device__ __host__ void skipahead(unsigned long long offset,
240  rocrand_state_sobol32* state)
241 {
242  return state->discard(offset);
243 }
244  // end of group rocranddevice
246 
247 #endif // ROCRAND_SOBOL32_H_
__forceinline__ __device__ __host__ unsigned int rocrand(rocrand_state_sobol32 *state)
Returns uniformly distributed random unsigned int value from [0; 2^32 - 1] range.
Definition: rocrand_sobol32.h:226
__forceinline__ __device__ __host__ void rocrand_init(const unsigned int *vectors, const unsigned int offset, rocrand_state_sobol32 *state)
Initialize SOBOL32 state.
Definition: rocrand_sobol32.h:207
__forceinline__ __device__ __host__ void skipahead(unsigned long long offset, rocrand_state_sobol32 *state)
Updates SOBOL32 state to skip ahead by offset elements.
Definition: rocrand_sobol32.h:239