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

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

API library: /build/reproducible-path/rocrand-6.4.1/library/include/rocrand/rocrand_lfsr113.h Source File
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
45namespace rocrand_device
46{
47namespace 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
72class lfsr113_engine
73{
74public:
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
165protected:
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
211protected:
212 lfsr113_state m_state;
213
214}; // lfsr113_engine class
215
216} // end namespace rocrand_device
217
224typedef 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