21 #ifndef ROCRAND_XORWOW_H_
22 #define ROCRAND_XORWOW_H_
25 #define FQUALIFIERS __forceinline__ __device__
28 #include "rocrand/rocrand_common.h"
29 #include "rocrand/rocrand_xorwow_precomputed.h"
39 #define ROCRAND_XORWOW_DEFAULT_SEED 0ULL
42 namespace rocrand_device {
46 void copy_vec(
unsigned int * dst,
const unsigned int * src)
48 for (
int i = 0; i < XORWOW_N; i++)
55 void mul_mat_vec_inplace(
const unsigned int * m,
unsigned int * v)
57 unsigned int r[XORWOW_N] = { 0 };
58 for (
int ij = 0; ij < XORWOW_N * XORWOW_M; ij++)
60 const int i = ij / XORWOW_M;
61 const int j = ij % XORWOW_M;
62 const unsigned int b = (v[i] & (1U << j)) ? 0xffffffff : 0x0;
63 for (
int k = 0; k < XORWOW_N; k++)
65 r[k] ^= b & m[i * XORWOW_M * XORWOW_N + j * XORWOW_N + k];
81 #ifndef ROCRAND_DETAIL_XORWOW_BM_NOT_IN_STATE
87 unsigned int boxmuller_float_state;
88 unsigned int boxmuller_double_state;
89 float boxmuller_float;
90 double boxmuller_double;
106 xorwow_engine(
const unsigned long long seed,
107 const unsigned long long subsequence,
108 const unsigned long long offset)
110 m_state.x[0] = 123456789U;
111 m_state.x[1] = 362436069U;
112 m_state.x[2] = 521288629U;
113 m_state.x[3] = 88675123U;
114 m_state.x[4] = 5783321U;
116 m_state.d = 6615241U;
119 const unsigned int s0 =
static_cast<unsigned int>(seed) ^ 0x2c7f967fU;
120 const unsigned int s1 =
static_cast<unsigned int>(seed >> 32) ^ 0xa03697cbU;
121 const unsigned int t0 = 1228688033U * s0;
122 const unsigned int t1 = 2073658381U * s1;
128 m_state.d += t1 + t0;
130 discard_subsequence(subsequence);
133 #ifndef ROCRAND_DETAIL_XORWOW_BM_NOT_IN_STATE
134 m_state.boxmuller_float_state = 0;
135 m_state.boxmuller_double_state = 0;
141 void discard(
unsigned long long offset)
143 #ifdef __HIP_DEVICE_COMPILE__
144 jump(offset, d_xorwow_jump_matrices);
146 jump(offset, h_xorwow_jump_matrices);
150 m_state.d +=
static_cast<unsigned int>(offset) * 362437;
156 void discard_subsequence(
unsigned long long subsequence)
159 #ifdef __HIP_DEVICE_COMPILE__
160 jump(subsequence, d_xorwow_sequence_jump_matrices);
162 jump(subsequence, h_xorwow_sequence_jump_matrices);
169 unsigned int operator()()
177 const unsigned int t = m_state.x[0] ^ (m_state.x[0] >> 2);
178 m_state.x[0] = m_state.x[1];
179 m_state.x[1] = m_state.x[2];
180 m_state.x[2] = m_state.x[3];
181 m_state.x[3] = m_state.x[4];
182 m_state.x[4] = (m_state.x[4] ^ (m_state.x[4] << 4)) ^ (t ^ (t << 1));
186 return m_state.d + m_state.x[4];
192 void jump(
unsigned long long v,
193 const unsigned int jump_matrices[XORWOW_JUMP_MATRICES][XORWOW_SIZE])
211 const unsigned int is =
static_cast<unsigned int>(v) & ((1 << XORWOW_JUMP_LOG2) - 1);
212 for (
unsigned int i = 0; i < is; i++)
214 detail::mul_mat_vec_inplace(jump_matrices[mi], m_state.x);
217 v >>= XORWOW_JUMP_LOG2;
223 xorwow_state m_state;
225 #ifndef ROCRAND_DETAIL_XORWOW_BM_NOT_IN_STATE
226 friend struct detail::engine_boxmuller_helper<xorwow_engine>;
239 typedef rocrand_device::xorwow_engine rocrand_state_xorwow;
255 const unsigned long long subsequence,
256 const unsigned long long offset,
257 rocrand_state_xorwow * state)
259 *state = rocrand_state_xorwow(seed, subsequence, offset);
275 unsigned int rocrand(rocrand_state_xorwow * state)
277 return state->next();
289 void skipahead(
unsigned long long offset, rocrand_state_xorwow * state)
291 return state->discard(offset);
306 return state->discard_subsequence(subsequence);
321 return state->discard_subsequence(sequence);
FQUALIFIERS void skipahead_subsequence(unsigned long long subsequence, rocrand_state_xorwow *state)
Updates XORWOW state to skip ahead by subsequence subsequences.
Definition: rocrand_xorwow.h:304
FQUALIFIERS unsigned int rocrand(rocrand_state_xorwow *state)
Returns uniformly distributed random unsigned int value from [0; 2^32 - 1] range.
Definition: rocrand_xorwow.h:275
FQUALIFIERS void rocrand_init(const unsigned long long seed, const unsigned long long subsequence, const unsigned long long offset, rocrand_state_xorwow *state)
Initialize XORWOW state.
Definition: rocrand_xorwow.h:254
FQUALIFIERS void skipahead(unsigned long long offset, rocrand_state_xorwow *state)
Updates XORWOW state to skip ahead by offset elements.
Definition: rocrand_xorwow.h:289
#define FQUALIFIERS
Shorthand for commonly used function qualifiers.
Definition: rocrand_uniform.h:31
#define ROCRAND_XORWOW_DEFAULT_SEED
Default seed for XORWOW PRNG.
Definition: rocrand_xorwow.h:39
FQUALIFIERS void skipahead_sequence(unsigned long long sequence, rocrand_state_xorwow *state)
Updates XORWOW state to skip ahead by sequence sequences.
Definition: rocrand_xorwow.h:319