Skip to main content
`Ctrl`+`K`

# /build/rocrand-D7zBeW/rocrand-7.2.4/library/include/rocrand/rocrand_sobol32.h Source File

API library: /build/rocrand-D7zBeW/rocrand-7.2.4/library/include/rocrand/rocrand_sobol32.h Source File
rocrand_sobol32.h
1// Copyright (c) 2017-2025 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 <hip/hip_runtime.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
192
194 typedef rocrand_device::sobol32_engine<false> rocrand_state_sobol32;
196
207__forceinline__ __device__ __host__
208 void rocrand_init(const unsigned int* vectors,
209 const unsigned int offset,
210 rocrand_state_sobol32* state)
211{
212 *state = rocrand_state_sobol32(vectors, offset);
213}
214
227__forceinline__ __device__ __host__
228 unsigned int rocrand(rocrand_state_sobol32* state)
229{
230 return state->next();
231}
232
241__forceinline__ __device__ __host__
242 void skipahead(unsigned long long offset, rocrand_state_sobol32* state)
243{
244 return state->discard(offset);
245}
246 // end of group rocranddevice
248
249#endif // ROCRAND_SOBOL32_H_
rocrand
__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:228
rocrand_init
__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:208
skipahead
__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:242