/build/reproducible-path/rocrand-6.4.1/library/include/rocrand/rocrand_mrg31k3p.h Source File

/build/reproducible-path/rocrand-6.4.1/library/include/rocrand/rocrand_mrg31k3p.h Source File#

API library: /build/reproducible-path/rocrand-6.4.1/library/include/rocrand/rocrand_mrg31k3p.h Source File
rocrand_mrg31k3p.h
1// Copyright (c) 2022-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_MRG31K3P_H_
22#define ROCRAND_MRG31K3P_H_
23
24#include "rocrand/rocrand_common.h"
25#include "rocrand/rocrand_mrg31k3p_precomputed.h"
26
27#define ROCRAND_MRG31K3P_M1 2147483647U // 2 ^ 31 - 1
28#define ROCRAND_MRG31K3P_M2 2147462579U // 2 ^ 31 - 21069
29#define ROCRAND_MRG31K3P_MASK12 511U // 2 ^ 9 - 1
30#define ROCRAND_MRG31K3P_MASK13 16777215U // 2 ^ 24 - 1
31#define ROCRAND_MRG31K3P_MASK21 65535U // 2 ^ 16 - 1
32#define ROCRAND_MRG31K3P_NORM_DOUBLE (4.656612875245796923e-10) // 1 / ROCRAND_MRG31K3P_M1
33#define ROCRAND_MRG31K3P_UINT32_NORM \
34 (2.000000001396983862) // UINT32_MAX / (ROCRAND_MRG31K3P_M1 - 1)
35
44#define ROCRAND_MRG31K3P_DEFAULT_SEED 12345ULL
// end of group rocranddevice
46
47namespace rocrand_device
48{
49
50class mrg31k3p_engine
51{
52public:
53 struct mrg31k3p_state
54 {
55 unsigned int x1[3];
56 unsigned int x2[3];
57
58#ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE
59 // The Box–Muller transform requires two inputs to convert uniformly
60 // distributed real values [0; 1] to normally distributed real values
61 // (with mean = 0, and stddev = 1). Often user wants only one
62 // normally distributed number, to save performance and random
63 // numbers the 2nd value is saved for future requests.
64 unsigned int boxmuller_float_state; // is there a float in boxmuller_float
65 unsigned int boxmuller_double_state; // is there a double in boxmuller_double
66 float boxmuller_float; // normally distributed float
67 double boxmuller_double; // normally distributed double
68#endif
69 };
70
71 __forceinline__ __device__ __host__ mrg31k3p_engine()
72 {
73 this->seed(ROCRAND_MRG31K3P_DEFAULT_SEED, 0, 0);
74 }
75
84 __forceinline__ __device__ __host__ mrg31k3p_engine(const unsigned long long seed,
85 const unsigned long long subsequence,
86 const unsigned long long offset)
87 {
88 this->seed(seed, subsequence, offset);
89 }
90
99 __forceinline__ __device__ __host__ void seed(unsigned long long seed_value,
100 const unsigned long long subsequence,
101 const unsigned long long offset)
102 {
103 if(seed_value == 0)
104 {
106 }
107 unsigned int x = static_cast<unsigned int>(seed_value ^ 0x55555555U);
108 unsigned int y = static_cast<unsigned int>((seed_value >> 32) ^ 0xAAAAAAAAU);
109 m_state.x1[0] = mod_mul_m1(x, seed_value);
110 m_state.x1[1] = mod_mul_m1(y, seed_value);
111 m_state.x1[2] = mod_mul_m1(x, seed_value);
112 m_state.x2[0] = mod_mul_m2(y, seed_value);
113 m_state.x2[1] = mod_mul_m2(x, seed_value);
114 m_state.x2[2] = mod_mul_m2(y, seed_value);
115 this->restart(subsequence, offset);
116 }
117
119 __forceinline__ __device__ __host__ void discard(unsigned long long offset)
120 {
121 this->discard_impl(offset);
122 }
123
126 __forceinline__ __device__ __host__ void discard_subsequence(unsigned long long subsequence)
127 {
128 this->discard_subsequence_impl(subsequence);
129 }
130
133 __forceinline__ __device__ __host__ void discard_sequence(unsigned long long sequence)
134 {
135 this->discard_sequence_impl(sequence);
136 }
137
138 __forceinline__ __device__ __host__ void restart(const unsigned long long subsequence,
139 const unsigned long long offset)
140 {
141#ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE
142 m_state.boxmuller_float_state = 0;
143 m_state.boxmuller_double_state = 0;
144#endif
145 this->discard_subsequence_impl(subsequence);
146 this->discard_impl(offset);
147 }
148
149 __forceinline__ __device__ __host__ unsigned int operator()()
150 {
151 return this->next();
152 }
153
154 // Returned value is in range [1, ROCRAND_MRG31K3P_M1].
155 __forceinline__ __device__ __host__ unsigned int next()
156 {
157 // First component
158 unsigned int tmp
159 = (((m_state.x1[1] & ROCRAND_MRG31K3P_MASK12) << 22) + (m_state.x1[1] >> 9))
160 + (((m_state.x1[2] & ROCRAND_MRG31K3P_MASK13) << 7) + (m_state.x1[2] >> 24));
161 tmp -= (tmp >= ROCRAND_MRG31K3P_M1) ? ROCRAND_MRG31K3P_M1 : 0;
162 tmp += m_state.x1[2];
163 tmp -= (tmp >= ROCRAND_MRG31K3P_M1) ? ROCRAND_MRG31K3P_M1 : 0;
164 m_state.x1[2] = m_state.x1[1];
165 m_state.x1[1] = m_state.x1[0];
166 m_state.x1[0] = tmp;
167
168 // Second component
169 tmp = (((m_state.x2[0] & ROCRAND_MRG31K3P_MASK21) << 15) + 21069 * (m_state.x2[0] >> 16));
170 tmp -= (tmp >= ROCRAND_MRG31K3P_M2) ? ROCRAND_MRG31K3P_M2 : 0;
171 tmp += ((m_state.x2[2] & ROCRAND_MRG31K3P_MASK21) << 15);
172 tmp -= (tmp >= ROCRAND_MRG31K3P_M2) ? ROCRAND_MRG31K3P_M2 : 0;
173 tmp += 21069 * (m_state.x2[2] >> 16);
174 tmp -= (tmp >= ROCRAND_MRG31K3P_M2) ? ROCRAND_MRG31K3P_M2 : 0;
175 tmp += m_state.x2[2];
176 tmp -= (tmp >= ROCRAND_MRG31K3P_M2) ? ROCRAND_MRG31K3P_M2 : 0;
177 m_state.x2[2] = m_state.x2[1];
178 m_state.x2[1] = m_state.x2[0];
179 m_state.x2[0] = tmp;
180
181 // Combination
182 return m_state.x1[0] - m_state.x2[0]
183 + (m_state.x1[0] <= m_state.x2[0] ? ROCRAND_MRG31K3P_M1 : 0);
184 }
185
186protected:
187 // Advances the internal state to skip \p offset numbers.
188 __forceinline__ __device__ __host__ void discard_impl(unsigned long long offset)
189 {
190 discard_state(offset);
191 }
192
193 // Advances the internal state to skip \p subsequence subsequences.
194 __forceinline__ __device__ __host__ void
195 discard_subsequence_impl(unsigned long long subsequence)
196 {
197 int i = 0;
198
199 while(subsequence > 0)
200 {
201 if(subsequence & 1)
202 {
203#if defined(__HIP_DEVICE_COMPILE__)
204 mod_mat_vec_m1(d_mrg31k3p_A1P72 + i, m_state.x1);
205 mod_mat_vec_m2(d_mrg31k3p_A2P72 + i, m_state.x2);
206#else
207 mod_mat_vec_m1(h_mrg31k3p_A1P72 + i, m_state.x1);
208 mod_mat_vec_m2(h_mrg31k3p_A2P72 + i, m_state.x2);
209#endif
210 }
211 subsequence >>= 1;
212 i += 9;
213 }
214 }
215
216 // Advances the internal state to skip \p sequences.
217 __forceinline__ __device__ __host__ void discard_sequence_impl(unsigned long long sequence)
218 {
219 int i = 0;
220
221 while(sequence > 0)
222 {
223 if(sequence & 1)
224 {
225#if defined(__HIP_DEVICE_COMPILE__)
226 mod_mat_vec_m1(d_mrg31k3p_A1P134 + i, m_state.x1);
227 mod_mat_vec_m2(d_mrg31k3p_A2P134 + i, m_state.x2);
228#else
229 mod_mat_vec_m1(h_mrg31k3p_A1P134 + i, m_state.x1);
230 mod_mat_vec_m2(h_mrg31k3p_A2P134 + i, m_state.x2);
231#endif
232 }
233 sequence >>= 1;
234 i += 9;
235 }
236 }
237
238 // Advances the internal state to skip \p offset numbers.
239 __forceinline__ __device__ __host__ void discard_state(unsigned long long offset)
240 {
241 int i = 0;
242
243 while(offset > 0)
244 {
245 if(offset & 1)
246 {
247#if defined(__HIP_DEVICE_COMPILE__)
248 mod_mat_vec_m1(d_mrg31k3p_A1 + i, m_state.x1);
249 mod_mat_vec_m2(d_mrg31k3p_A2 + i, m_state.x2);
250#else
251 mod_mat_vec_m1(h_mrg31k3p_A1 + i, m_state.x1);
252 mod_mat_vec_m2(h_mrg31k3p_A2 + i, m_state.x2);
253#endif
254 }
255 offset >>= 1;
256 i += 9;
257 }
258 }
259
260 // Advances the internal state to the next state.
261 __forceinline__ __device__ __host__ void discard_state()
262 {
263 discard_state(1);
264 }
265
266private:
267 __forceinline__ __device__ __host__ static void mod_mat_vec_m1(const unsigned int* A,
268 unsigned int* s)
269 {
270 unsigned long long x[3] = {s[0], s[1], s[2]};
271
272 s[0] = mod_m1(mod_m1(A[0] * x[0]) + mod_m1(A[1] * x[1]) + mod_m1(A[2] * x[2]));
273
274 s[1] = mod_m1(mod_m1(A[3] * x[0]) + mod_m1(A[4] * x[1]) + mod_m1(A[5] * x[2]));
275
276 s[2] = mod_m1(mod_m1(A[6] * x[0]) + mod_m1(A[7] * x[1]) + mod_m1(A[8] * x[2]));
277 }
278
279 __forceinline__ __device__ __host__ static void mod_mat_vec_m2(const unsigned int* A,
280 unsigned int* s)
281 {
282 unsigned long long x[3] = {s[0], s[1], s[2]};
283
284 s[0] = mod_m2(mod_m2(A[0] * x[0]) + mod_m2(A[1] * x[1]) + mod_m2(A[2] * x[2]));
285
286 s[1] = mod_m2(mod_m2(A[3] * x[0]) + mod_m2(A[4] * x[1]) + mod_m2(A[5] * x[2]));
287
288 s[2] = mod_m2(mod_m2(A[6] * x[0]) + mod_m2(A[7] * x[1]) + mod_m2(A[8] * x[2]));
289 }
290
291 __forceinline__ __device__ __host__ static unsigned long long mod_mul_m1(unsigned int i,
292 unsigned long long j)
293 {
294 return mod_m1(i * j);
295 }
296
297 __forceinline__ __device__ __host__ static unsigned long long mod_m1(unsigned long long p)
298 {
299 return p % ROCRAND_MRG31K3P_M1;
300 }
301
302 __forceinline__ __device__ __host__ static unsigned long long mod_mul_m2(unsigned int i,
303 unsigned long long j)
304 {
305 return mod_m2(i * j);
306 }
307
308 __forceinline__ __device__ __host__ static unsigned long long mod_m2(unsigned long long p)
309 {
310 return p % ROCRAND_MRG31K3P_M2;
311 }
312
313protected:
314 // State
315 mrg31k3p_state m_state;
316
317#ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE
318 friend struct detail::engine_boxmuller_helper<mrg31k3p_engine>;
319#endif
320}; // mrg31k3p_engine class
321
322} // end namespace rocrand_device
323
330typedef rocrand_device::mrg31k3p_engine rocrand_state_mrg31k3p;
332
344__forceinline__ __device__ __host__ void rocrand_init(const unsigned long long seed,
345 const unsigned long long subsequence,
346 const unsigned long long offset,
347 rocrand_state_mrg31k3p* state)
348{
349 *state = rocrand_state_mrg31k3p(seed, subsequence, offset);
350}
351
364__forceinline__ __device__ __host__ unsigned int rocrand(rocrand_state_mrg31k3p* state)
365{
366 // next() in [1, ROCRAND_MRG31K3P_M1]
367 return static_cast<unsigned int>((state->next() - 1) * ROCRAND_MRG31K3P_UINT32_NORM);
368}
369
378__forceinline__ __device__ __host__ void skipahead(unsigned long long offset,
379 rocrand_state_mrg31k3p* state)
380{
381 return state->discard(offset);
382}
383
393__forceinline__ __device__ __host__ void skipahead_subsequence(unsigned long long subsequence,
394 rocrand_state_mrg31k3p* state)
395{
396 return state->discard_subsequence(subsequence);
397}
398
408__forceinline__ __device__ __host__ void skipahead_sequence(unsigned long long sequence,
409 rocrand_state_mrg31k3p* state)
410{
411 return state->discard_sequence(sequence);
412}
413
// end of group rocranddevice
415
416#endif // ROCRAND_MRG31K3P_H_
#define ROCRAND_MRG31K3P_DEFAULT_SEED
Default seed for MRG31K3P PRNG.
Definition rocrand_mrg31k3p.h:44
__forceinline__ __device__ __host__ unsigned int rocrand(rocrand_state_mrg31k3p *state)
Returns uniformly distributed random unsigned int value from [0; 2^32 - 1] range.
Definition rocrand_mrg31k3p.h:364
__forceinline__ __device__ __host__ void skipahead_subsequence(unsigned long long subsequence, rocrand_state_mrg31k3p *state)
Updates MRG31K3P state to skip ahead by subsequence subsequences.
Definition rocrand_mrg31k3p.h:393
__forceinline__ __device__ __host__ void skipahead(unsigned long long offset, rocrand_state_mrg31k3p *state)
Updates MRG31K3P state to skip ahead by offset elements.
Definition rocrand_mrg31k3p.h:378
__forceinline__ __device__ __host__ void skipahead_sequence(unsigned long long sequence, rocrand_state_mrg31k3p *state)
Updates MRG31K3P state to skip ahead by sequence sequences.
Definition rocrand_mrg31k3p.h:408
__forceinline__ __device__ __host__ void rocrand_init(const unsigned long long seed, const unsigned long long subsequence, const unsigned long long offset, rocrand_state_mrg31k3p *state)
Initializes MRG31K3P state.
Definition rocrand_mrg31k3p.h:344