53 #ifndef ROCRAND_PHILOX4X32_10_H_
54 #define ROCRAND_PHILOX4X32_10_H_
56 #include "rocrand/rocrand_common.h"
58 #include <hip/hip_runtime.h>
62 #define ROCRAND_PHILOX_M4x32_0 0xD2511F53U
63 #define ROCRAND_PHILOX_M4x32_1 0xCD9E8D57U
64 #define ROCRAND_PHILOX_W32_0 0x9E3779B9U
65 #define ROCRAND_PHILOX_W32_1 0xBB67AE85U
67 #if defined(__HIP_DEVICE_COMPILE__) && defined(__AMDGCN__)
68 #define ROCRAND_BUILTIN_ADDC(a, b, carry_in, carry_out) \
69 __builtin_addc(a, b, carry_in, carry_out)
71 #define ROCRAND_BUILTIN_ADDC(a, b, carry_in, carry_out) \
75 unsigned int sum = (a) + (b) + (carry_in); \
76 *carry_out = (sum < (a) || ((carry_in) && sum <= (b))) ? 1 : 0; \
89 #define ROCRAND_PHILOX4x32_DEFAULT_SEED 0xdeadbeefdeadbeefULL
92 namespace rocrand_device
95 class philox4x32_10_engine
98 struct philox4x32_10_state
103 unsigned int substate;
105 #ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE
111 float boxmuller_float;
112 double boxmuller_double;
116 __forceinline__ __device__ __host__
117 philox4x32_10_engine()
127 __forceinline__ __device__ __host__
128 philox4x32_10_engine(
const unsigned long long seed,
129 const unsigned long long subsequence,
130 const unsigned long long offset)
132 this->seed(seed, subsequence, offset);
140 __forceinline__ __device__ __host__
141 void seed(
unsigned long long seed_value,
142 const unsigned long long subsequence,
143 const unsigned long long offset)
145 m_state.key.x =
static_cast<unsigned int>(seed_value);
146 m_state.key.y =
static_cast<unsigned int>(seed_value >> 32);
147 this->restart(subsequence, offset);
151 __forceinline__ __device__ __host__
152 void discard(
unsigned long long offset)
154 this->discard_impl(offset);
155 this->m_state.result = this->ten_rounds(m_state.counter, m_state.key);
162 __forceinline__ __device__ __host__
163 void discard_subsequence(
unsigned long long subsequence)
165 this->discard_subsequence_impl(subsequence);
166 m_state.result = this->ten_rounds(m_state.counter, m_state.key);
169 __forceinline__ __device__ __host__
170 void restart(
const unsigned long long subsequence,
const unsigned long long offset)
172 m_state.counter = {0, 0, 0, 0};
173 m_state.result = {0, 0, 0, 0};
174 m_state.substate = 0;
175 #ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE
176 m_state.boxmuller_float = ROCRAND_NAN_FLOAT;
177 m_state.boxmuller_double = ROCRAND_NAN_DOUBLE;
179 this->discard_subsequence_impl(subsequence);
180 this->discard_impl(offset);
181 m_state.result = this->ten_rounds(m_state.counter, m_state.key);
184 __forceinline__ __device__ __host__
185 unsigned int operator()()
190 __forceinline__ __device__ __host__
193 unsigned int s = m_state.substate;
195 = {m_state.result.x, m_state.result.y, m_state.result.z, m_state.result.w};
196 unsigned int ret = vs[s];
202 this->discard_state();
203 m_state.result = this->ten_rounds(m_state.counter, m_state.key);
205 m_state.substate = s;
209 __forceinline__ __device__ __host__
212 uint4 ret = m_state.result;
213 this->discard_state();
214 m_state.result = this->ten_rounds(m_state.counter, m_state.key);
215 return this->interleave(ret, m_state.result);
221 __forceinline__ __device__ __host__
222 void discard_impl(
unsigned long long offset)
225 m_state.substate += offset & 3;
226 unsigned long long counter_offset = offset / 4;
227 counter_offset += m_state.substate < 4 ? 0 : 1;
228 m_state.substate += m_state.substate < 4 ? 0 : -4;
230 this->discard_state(counter_offset);
234 __forceinline__ __device__ __host__
235 void discard_subsequence_impl(
unsigned long long subsequence)
237 const unsigned int lo =
static_cast<unsigned int>(subsequence);
238 const unsigned int hi =
static_cast<unsigned int>(subsequence >> 32);
239 unsigned int carry = 0;
241 m_state.counter.z = ROCRAND_BUILTIN_ADDC(m_state.counter.z, lo, 0, &carry);
242 m_state.counter.w = ROCRAND_BUILTIN_ADDC(m_state.counter.w, hi, carry, &carry);
247 __forceinline__ __device__ __host__
248 void discard_state(
unsigned long long offset)
250 const unsigned int lo =
static_cast<unsigned int>(offset);
251 const unsigned int hi =
static_cast<unsigned int>(offset >> 32);
252 unsigned int carry = 0;
254 m_state.counter.x = ROCRAND_BUILTIN_ADDC(m_state.counter.x, lo, 0, &carry);
255 m_state.counter.y = ROCRAND_BUILTIN_ADDC(m_state.counter.y, hi, carry, &carry);
256 m_state.counter.z = ROCRAND_BUILTIN_ADDC(m_state.counter.z, 0, carry, &carry);
257 m_state.counter.w = ROCRAND_BUILTIN_ADDC(m_state.counter.w, 0, carry, &carry);
262 __forceinline__ __device__ __host__
265 m_state.counter = this->bump_counter(m_state.counter);
268 __forceinline__ __device__ __host__
269 static uint4 bump_counter(uint4 counter)
271 unsigned int carry = 0;
272 counter.x = ROCRAND_BUILTIN_ADDC(counter.x, 1, 0, &carry);
273 counter.y = ROCRAND_BUILTIN_ADDC(counter.y, 0, carry, &carry);
274 counter.z = ROCRAND_BUILTIN_ADDC(counter.z, 0, carry, &carry);
275 counter.w = ROCRAND_BUILTIN_ADDC(counter.w, 0, carry, &carry);
279 __forceinline__ __device__ __host__
280 uint4 interleave(
const uint4 prev,
const uint4 next)
const
282 switch(m_state.substate)
285 case 1:
return uint4{prev.y, prev.z, prev.w, next.x};
286 case 2:
return uint4{prev.z, prev.w, next.x, next.y};
287 case 3:
return uint4{prev.w, next.x, next.y, next.z};
289 __builtin_unreachable();
293 __forceinline__ __device__ __host__
294 uint4 ten_rounds(uint4 counter, uint2 key)
296 counter = this->single_round(counter, key);
297 key = this->bumpkey(key);
298 counter = this->single_round(counter, key);
299 key = this->bumpkey(key);
300 counter = this->single_round(counter, key);
301 key = this->bumpkey(key);
302 counter = this->single_round(counter, key);
303 key = this->bumpkey(key);
304 counter = this->single_round(counter, key);
305 key = this->bumpkey(key);
306 counter = this->single_round(counter, key);
307 key = this->bumpkey(key);
308 counter = this->single_round(counter, key);
309 key = this->bumpkey(key);
310 counter = this->single_round(counter, key);
311 key = this->bumpkey(key);
312 counter = this->single_round(counter, key);
313 key = this->bumpkey(key);
314 return this->single_round(counter, key);
319 __forceinline__ __device__ __host__
320 static uint4 single_round(uint4 counter, uint2 key)
323 unsigned long long mul0 = detail::mul_u64_u32(ROCRAND_PHILOX_M4x32_0, counter.x);
324 unsigned int hi0 =
static_cast<unsigned int>(mul0 >> 32);
325 unsigned int lo0 =
static_cast<unsigned int>(mul0);
326 unsigned long long mul1 = detail::mul_u64_u32(ROCRAND_PHILOX_M4x32_1, counter.z);
327 unsigned int hi1 =
static_cast<unsigned int>(mul1 >> 32);
328 unsigned int lo1 =
static_cast<unsigned int>(mul1);
329 return uint4{hi1 ^ counter.y ^ key.x, lo1, hi0 ^ counter.w ^ key.y, lo0};
332 __forceinline__ __device__ __host__
333 static uint2 bumpkey(uint2 key)
335 key.x += ROCRAND_PHILOX_W32_0;
336 key.y += ROCRAND_PHILOX_W32_1;
342 philox4x32_10_state m_state;
344 #ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE
345 friend struct detail::engine_boxmuller_helper<philox4x32_10_engine>;
358 typedef rocrand_device::philox4x32_10_engine rocrand_state_philox4x32_10;
372 __forceinline__ __device__ __host__
374 const unsigned long long subsequence,
375 const unsigned long long offset,
376 rocrand_state_philox4x32_10* state)
378 *state = rocrand_state_philox4x32_10(seed, subsequence, offset);
393 __forceinline__ __device__ __host__
394 unsigned int rocrand(rocrand_state_philox4x32_10* state)
396 return state->next();
411 __forceinline__ __device__ __host__
414 return state->next4();
425 __forceinline__ __device__ __host__
426 void skipahead(
unsigned long long offset, rocrand_state_philox4x32_10* state)
428 return state->discard(offset);
440 __forceinline__ __device__ __host__
443 return state->discard_subsequence(subsequence);
455 __forceinline__ __device__ __host__
458 return state->discard_subsequence(sequence);
__forceinline__ __device__ __host__ void skipahead(unsigned long long offset, rocrand_state_philox4x32_10 *state)
Updates Philox state to skip ahead by offset elements.
Definition: rocrand_philox4x32_10.h:426
__forceinline__ __device__ __host__ void rocrand_init(const unsigned long long seed, const unsigned long long subsequence, const unsigned long long offset, rocrand_state_philox4x32_10 *state)
Initializes Philox state.
Definition: rocrand_philox4x32_10.h:373
#define ROCRAND_PHILOX4x32_DEFAULT_SEED
Default seed for PHILOX4x32 PRNG.
Definition: rocrand_philox4x32_10.h:89
__forceinline__ __device__ __host__ uint4 rocrand4(rocrand_state_philox4x32_10 *state)
Returns four uniformly distributed random unsigned int values from [0; 2^32 - 1] range.
Definition: rocrand_philox4x32_10.h:412
__forceinline__ __device__ __host__ void skipahead_sequence(unsigned long long sequence, rocrand_state_philox4x32_10 *state)
Updates Philox state to skip ahead by sequence sequences.
Definition: rocrand_philox4x32_10.h:456
__forceinline__ __device__ __host__ unsigned int rocrand(rocrand_state_philox4x32_10 *state)
Returns uniformly distributed random unsigned int value from [0; 2^32 - 1] range.
Definition: rocrand_philox4x32_10.h:394
__forceinline__ __device__ __host__ void skipahead_subsequence(unsigned long long subsequence, rocrand_state_philox4x32_10 *state)
Updates Philox state to skip ahead by subsequence subsequences.
Definition: rocrand_philox4x32_10.h:441