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

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-rocrand/checkouts/docs-6.2.1/library/include/rocrand/rocrand_lfsr113.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_lfsr113.h Source File
API library
rocrand_lfsr113.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_LFSR113_H_
22 #define ROCRAND_LFSR113_H_
23 
24 #include "rocrand/rocrand_common.h"
25 #include "rocrand/rocrand_lfsr113_precomputed.h"
26 
33 #define ROCRAND_LFSR113_DEFAULT_SEED_X 2
36 #define ROCRAND_LFSR113_DEFAULT_SEED_Y 8
39 #define ROCRAND_LFSR113_DEFAULT_SEED_Z 16
42 #define ROCRAND_LFSR113_DEFAULT_SEED_W 128 // end of group rocranddevice
44 
45 namespace rocrand_device
46 {
47 namespace detail
48 {
49 
50 __forceinline__ __device__ __host__ void mul_mat_vec_inplace(const unsigned int* m, uint4* z)
51 {
52  unsigned int v[4] = {z->x, z->y, z->z, z->w};
53  unsigned int r[LFSR113_N] = {0};
54  for(int ij = 0; ij < LFSR113_N * LFSR113_M; ij++)
55  {
56  const int i = ij / LFSR113_M;
57  const int j = ij % LFSR113_M;
58  const unsigned int b = (v[i] & (1U << j)) ? 0xffffffff : 0x0;
59  for(int k = 0; k < LFSR113_N; k++)
60  {
61  r[k] ^= b & m[i * LFSR113_M * LFSR113_N + j * LFSR113_N + k];
62  }
63  }
64  // Copy result into z
65  z->x = r[0];
66  z->y = r[1];
67  z->z = r[2];
68  z->w = r[3];
69 }
70 } // end namespace detail
71 
72 class lfsr113_engine
73 {
74 public:
75  struct lfsr113_state
76  {
77  uint4 z;
78  uint4 subsequence;
79  };
80 
86  __forceinline__ __device__ __host__ lfsr113_engine(const uint4 seed
91  const unsigned int subsequence = 0,
92  const unsigned long long offset = 0)
93  {
94  this->seed(seed, subsequence, offset);
95  }
96 
102  __forceinline__ __device__ __host__ void seed(uint4 seed_value,
103  const unsigned long long subsequence,
104  const unsigned long long offset = 0)
105  {
106  m_state.subsequence = seed_value;
107 
108  reset_start_subsequence();
109  discard_subsequence(subsequence);
110  discard(offset);
111  }
112 
114  __forceinline__ __device__ __host__ void discard()
115  {
116  discard_state();
117  }
118 
120  __forceinline__ __device__ __host__ void discard(unsigned long long offset)
121  {
122 #ifdef __HIP_DEVICE_COMPILE__
123  jump(offset, d_lfsr113_jump_matrices);
124 #else
125  jump(offset, h_lfsr113_jump_matrices);
126 #endif
127  }
128 
131  __forceinline__ __device__ __host__ void discard_subsequence(unsigned int subsequence)
132  {
133 // Discard n * 2^55 samples
134 #ifdef __HIP_DEVICE_COMPILE__
135  jump(subsequence, d_lfsr113_sequence_jump_matrices);
136 #else
137  jump(subsequence, h_lfsr113_sequence_jump_matrices);
138 #endif
139  }
140 
141  __forceinline__ __device__ __host__ unsigned int operator()()
142  {
143  return next();
144  }
145 
146  __forceinline__ __device__ __host__ unsigned int next()
147  {
148  unsigned int b;
149 
150  b = (((m_state.z.x << 6) ^ m_state.z.x) >> 13);
151  m_state.z.x = (((m_state.z.x & 4294967294U) << 18) ^ b);
152 
153  b = (((m_state.z.y << 2) ^ m_state.z.y) >> 27);
154  m_state.z.y = (((m_state.z.y & 4294967288U) << 2) ^ b);
155 
156  b = (((m_state.z.z << 13) ^ m_state.z.z) >> 21);
157  m_state.z.z = (((m_state.z.z & 4294967280U) << 7) ^ b);
158 
159  b = (((m_state.z.w << 3) ^ m_state.z.w) >> 12);
160  m_state.z.w = (((m_state.z.w & 4294967168U) << 13) ^ b);
161 
162  return (m_state.z.x ^ m_state.z.y ^ m_state.z.z ^ m_state.z.w);
163  }
164 
165 protected:
167  __forceinline__ __device__ __host__ void reset_start_subsequence()
168  {
169  m_state.z.x = m_state.subsequence.x;
170  m_state.z.y = m_state.subsequence.y;
171  m_state.z.z = m_state.subsequence.z;
172  m_state.z.w = m_state.subsequence.w;
173  }
174 
175  // Advances the internal state to the next state.
176  __forceinline__ __device__ __host__ void discard_state()
177  {
178  this->next();
179  }
180 
181  __forceinline__ __device__ __host__ void
182  jump(unsigned long long v,
183  const unsigned int (&jump_matrices)[LFSR113_JUMP_MATRICES][LFSR113_SIZE])
184  {
185  // x~(n + v) = (A^v mod m)x~n mod m
186  // The matrix (A^v mod m) can be precomputed for selected values of v.
187  //
188  // For LFSR113_JUMP_LOG2 = 2
189  // lfsr113_jump_matrices contains precomputed matrices:
190  // A^1, A^4, A^16...
191  //
192  // For LFSR113_JUMP_LOG2 = 2 and LFSR113_SEQUENCE_JUMP_LOG2 = 55
193  // lfsr113_sequence_jump_matrices contains precomputed matrices:
194  // A^(1 * 2^55), A^(4 * 2^55), A^(16 * 2^55)...
195  //
196  // Intermediate powers can be calculated as multiplication of the powers above.
197 
198  unsigned int mi = 0;
199  while(v > 0)
200  {
201  const unsigned int is = static_cast<unsigned int>(v) & ((1 << LFSR113_JUMP_LOG2) - 1);
202  for(unsigned int i = 0; i < is; i++)
203  {
204  detail::mul_mat_vec_inplace(jump_matrices[mi], &m_state.z);
205  }
206  mi++;
207  v >>= LFSR113_JUMP_LOG2;
208  }
209  }
210 
211 protected:
212  lfsr113_state m_state;
213 
214 }; // lfsr113_engine class
215 
216 } // end namespace rocrand_device
217 
224 typedef rocrand_device::lfsr113_engine rocrand_state_lfsr113;
226 
237 __forceinline__ __device__ __host__ void
238  rocrand_init(const uint4 seed, const unsigned int subsequence, rocrand_state_lfsr113* state)
239 {
240  *state = rocrand_state_lfsr113(seed, subsequence);
241 }
242 
254 __forceinline__ __device__ __host__ void rocrand_init(const uint4 seed,
255  const unsigned int subsequence,
256  const unsigned long long offset,
257  rocrand_state_lfsr113* state)
258 {
259  *state = rocrand_state_lfsr113(seed, subsequence, offset);
260 }
261 
274 __forceinline__ __device__ __host__ unsigned int rocrand(rocrand_state_lfsr113* state)
275 {
276  return state->next();
277 }
278 
287 __forceinline__ __device__ __host__ void skipahead(unsigned long long offset,
288  rocrand_state_lfsr113* state)
289 {
290  return state->discard(offset);
291 }
292 
302 __forceinline__ __device__ __host__ void skipahead_subsequence(unsigned int subsequence,
303  rocrand_state_lfsr113* state)
304 {
305  return state->discard_subsequence(subsequence);
306 }
307 
317 __forceinline__ __device__ __host__ void skipahead_sequence(unsigned int sequence,
318  rocrand_state_lfsr113* state)
319 {
320  return state->discard_subsequence(sequence);
321 }
322  // end of group rocranddevice
324 
325 #endif // ROCRAND_LFSR113_H_
__forceinline__ __device__ __host__ void rocrand_init(const uint4 seed, const unsigned int subsequence, rocrand_state_lfsr113 *state)
Initializes LFSR113 state.
Definition: rocrand_lfsr113.h:238
__forceinline__ __device__ __host__ void skipahead(unsigned long long offset, rocrand_state_lfsr113 *state)
Updates LFSR113 state to skip ahead by offset elements.
Definition: rocrand_lfsr113.h:287
#define ROCRAND_LFSR113_DEFAULT_SEED_Y
Default Y seed for LFSR113 PRNG.
Definition: rocrand_lfsr113.h:36
__forceinline__ __device__ __host__ void skipahead_sequence(unsigned int sequence, rocrand_state_lfsr113 *state)
Updates LFSR113 state to skip ahead by sequence sequences.
Definition: rocrand_lfsr113.h:317
#define ROCRAND_LFSR113_DEFAULT_SEED_W
Default W seed for LFSR113 PRNG.
Definition: rocrand_lfsr113.h:42
__forceinline__ __device__ __host__ void skipahead_subsequence(unsigned int subsequence, rocrand_state_lfsr113 *state)
Updates LFSR113 state to skip ahead by subsequence subsequences.
Definition: rocrand_lfsr113.h:302
#define ROCRAND_LFSR113_DEFAULT_SEED_Z
Default Z seed for LFSR113 PRNG.
Definition: rocrand_lfsr113.h:39
#define ROCRAND_LFSR113_DEFAULT_SEED_X
Default X seed for LFSR113 PRNG.
Definition: rocrand_lfsr113.h:33
__forceinline__ __device__ __host__ unsigned int rocrand(rocrand_state_lfsr113 *state)
Returns uniformly distributed random unsigned int value from [0; 2^32 - 1] range.
Definition: rocrand_lfsr113.h:274