26#ifndef ROCRAND_UNIFORM_H_
27#define ROCRAND_UNIFORM_H_
29#include "rocrand/rocrand_lfsr113.h"
30#include "rocrand/rocrand_mrg31k3p.h"
31#include "rocrand/rocrand_mrg32k3a.h"
32#include "rocrand/rocrand_mtgp32.h"
33#include "rocrand/rocrand_philox4x32_10.h"
34#include "rocrand/rocrand_scrambled_sobol32.h"
35#include "rocrand/rocrand_scrambled_sobol64.h"
36#include "rocrand/rocrand_sobol32.h"
37#include "rocrand/rocrand_sobol64.h"
38#include "rocrand/rocrand_threefry2x32_20.h"
39#include "rocrand/rocrand_threefry2x64_20.h"
40#include "rocrand/rocrand_threefry4x32_20.h"
41#include "rocrand/rocrand_threefry4x64_20.h"
42#include "rocrand/rocrand_xorwow.h"
44#include "rocrand/rocrand_common.h"
46namespace rocrand_device {
55union two_uints_to_ulong
57 two_uints uint2_value;
58 unsigned long long int ulong_value;
63__forceinline__ __device__ __host__
float uniform_distribution(
unsigned int v)
65 return ROCRAND_2POW32_INV + (v * ROCRAND_2POW32_INV);
70__forceinline__ __device__ __host__
float uniform_distribution(
unsigned long long int v)
72 return ROCRAND_2POW32_INV + (v >> 32) * ROCRAND_2POW32_INV;
75__forceinline__ __device__ __host__ float4 uniform_distribution4(uint4 v)
78 ROCRAND_2POW32_INV + (v.x * ROCRAND_2POW32_INV),
79 ROCRAND_2POW32_INV + (v.y * ROCRAND_2POW32_INV),
80 ROCRAND_2POW32_INV + (v.z * ROCRAND_2POW32_INV),
81 ROCRAND_2POW32_INV + (v.w * ROCRAND_2POW32_INV)
85__forceinline__ __device__ __host__ float4 uniform_distribution4(ulonglong4 v)
87 return float4{ROCRAND_2POW64_INV + (v.x * ROCRAND_2POW64_INV),
88 ROCRAND_2POW64_INV + (v.y * ROCRAND_2POW64_INV),
89 ROCRAND_2POW64_INV + (v.z * ROCRAND_2POW64_INV),
90 ROCRAND_2POW64_INV + (v.w * ROCRAND_2POW64_INV)};
95__forceinline__ __device__ __host__
double uniform_distribution_double(
unsigned int v)
97 return ROCRAND_2POW32_INV_DOUBLE + (v * ROCRAND_2POW32_INV_DOUBLE);
100__forceinline__ __device__ __host__
double uniform_distribution_double(
unsigned int v1,
103 two_uints_to_ulong v;
104 v.uint2_value.x = v1;
105 v.uint2_value.y = (v2 >> 11);
106 return ROCRAND_2POW53_INV_DOUBLE + (v.ulong_value * ROCRAND_2POW53_INV_DOUBLE);
109__forceinline__ __device__ __host__
double uniform_distribution_double(
unsigned long long int v)
111 return ROCRAND_2POW53_INV_DOUBLE + (
114 (v >> 11) * ROCRAND_2POW53_INV_DOUBLE
118__forceinline__ __device__ __host__ double2 uniform_distribution_double2(uint4 v)
121 uniform_distribution_double(v.x, v.y),
122 uniform_distribution_double(v.z, v.w)
126__forceinline__ __device__ __host__ double4 uniform_distribution_double4(uint4 v1, uint4 v2)
129 uniform_distribution_double(v1.x, v1.y),
130 uniform_distribution_double(v1.z, v1.w),
131 uniform_distribution_double(v2.x, v2.y),
132 uniform_distribution_double(v2.z, v2.w)
136__forceinline__ __device__ __host__ double2 uniform_distribution_double2(ulonglong2 v)
138 return double2{uniform_distribution_double(v.x), uniform_distribution_double(v.y)};
141__forceinline__ __device__ __host__ double2 uniform_distribution_double2(ulonglong4 v)
143 return double2{uniform_distribution_double(v.x), uniform_distribution_double(v.y)};
146__forceinline__ __device__ __host__ double4 uniform_distribution_double4(ulonglong4 v)
148 return double4{uniform_distribution_double(v.x),
149 uniform_distribution_double(v.z),
150 uniform_distribution_double(v.x),
151 uniform_distribution_double(v.z)};
154__forceinline__ __device__ __host__ __half uniform_distribution_half(
unsigned short v)
156 return __float2half(ROCRAND_2POW16_INV + (v * ROCRAND_2POW16_INV));
161template<
typename state_type>
162__forceinline__ __device__ __host__
unsigned int mrg_uniform_distribution_uint(
unsigned int v)
166__forceinline__ __device__ __host__
unsigned int
167 mrg_uniform_distribution_uint<rocrand_state_mrg31k3p>(
unsigned int v)
169 return static_cast<unsigned int>((v - 1) * ROCRAND_MRG31K3P_UINT32_NORM);
173__forceinline__ __device__ __host__
unsigned int
174 mrg_uniform_distribution_uint<rocrand_state_mrg32k3a>(
unsigned int v)
176 return static_cast<unsigned int>((v - 1) * ROCRAND_MRG32K3A_UINT_NORM);
181template<
typename state_type>
182__forceinline__ __device__ __host__
float mrg_uniform_distribution(
unsigned int v) =
delete;
185__forceinline__ __device__ __host__
float
186 mrg_uniform_distribution<rocrand_state_mrg31k3p>(
unsigned int v)
188 double ret =
static_cast<double>(v) * ROCRAND_MRG31K3P_NORM_DOUBLE;
189 return static_cast<float>(ret);
193__forceinline__ __device__ __host__
float
194 mrg_uniform_distribution<rocrand_state_mrg32k3a>(
unsigned int v)
196 double ret =
static_cast<double>(v) * ROCRAND_MRG32K3A_NORM_DOUBLE;
197 return static_cast<float>(ret);
202template<
typename state_type>
203__forceinline__ __device__ __host__
double mrg_uniform_distribution_double(
unsigned int v) =
delete;
206__forceinline__ __device__ __host__
double
207 mrg_uniform_distribution_double<rocrand_state_mrg31k3p>(
unsigned int v)
209 double ret =
static_cast<double>(v) * ROCRAND_MRG31K3P_NORM_DOUBLE;
214__forceinline__ __device__ __host__
double
215 mrg_uniform_distribution_double<rocrand_state_mrg32k3a>(
unsigned int v)
217 double ret =
static_cast<double>(v) * ROCRAND_MRG32K3A_NORM_DOUBLE;
236__forceinline__ __device__ __host__
float rocrand_uniform(rocrand_state_philox4x32_10* state)
238 return rocrand_device::detail::uniform_distribution(
rocrand(state));
253__forceinline__ __device__ __host__ float2
rocrand_uniform2(rocrand_state_philox4x32_10* state)
259 rocrand_device::detail::uniform_distribution(state1),
260 rocrand_device::detail::uniform_distribution(state2)
276__forceinline__ __device__ __host__ float4
rocrand_uniform4(rocrand_state_philox4x32_10* state)
278 return rocrand_device::detail::uniform_distribution4(
rocrand4(state));
293__forceinline__ __device__ __host__
double
299 return rocrand_device::detail::uniform_distribution_double(state1, state2);
314__forceinline__ __device__ __host__ double2
317 return rocrand_device::detail::uniform_distribution_double2(
rocrand4(state));
332__forceinline__ __device__ __host__ double4
335 return rocrand_device::detail::uniform_distribution_double4(
rocrand4(state),
rocrand4(state));
350__forceinline__ __device__ __host__
float rocrand_uniform(rocrand_state_mrg31k3p* state)
352 return rocrand_device::detail::mrg_uniform_distribution<rocrand_state_mrg31k3p>(state->next());
372 return rocrand_device::detail::mrg_uniform_distribution_double<rocrand_state_mrg31k3p>(
388__forceinline__ __device__ __host__
float rocrand_uniform(rocrand_state_mrg32k3a* state)
390 return rocrand_device::detail::mrg_uniform_distribution<rocrand_state_mrg32k3a>(state->next());
410 return rocrand_device::detail::mrg_uniform_distribution_double<rocrand_state_mrg32k3a>(
428 return rocrand_device::detail::uniform_distribution(
rocrand(state));
448 return rocrand_device::detail::uniform_distribution_double(state1, state2);
465 return rocrand_device::detail::uniform_distribution(
rocrand(state));
485 return rocrand_device::detail::uniform_distribution_double(
rocrand(state));
500__forceinline__ __device__ __host__
float rocrand_uniform(rocrand_state_sobol32* state)
502 return rocrand_device::detail::uniform_distribution(
rocrand(state));
522 return rocrand_device::detail::uniform_distribution_double(
rocrand(state));
537__forceinline__ __device__ __host__
float rocrand_uniform(rocrand_state_scrambled_sobol32* state)
539 return rocrand_device::detail::uniform_distribution(
rocrand(state));
557__forceinline__ __device__ __host__
double
560 return rocrand_device::detail::uniform_distribution_double(
rocrand(state));
575__forceinline__ __device__ __host__
float rocrand_uniform(rocrand_state_sobol64* state)
577 return rocrand_device::detail::uniform_distribution(
rocrand(state));
594 return rocrand_device::detail::uniform_distribution_double(
rocrand(state));
609__forceinline__ __device__ __host__
float rocrand_uniform(rocrand_state_scrambled_sobol64* state)
611 return rocrand_device::detail::uniform_distribution(
rocrand(state));
626__forceinline__ __device__ __host__
double
629 return rocrand_device::detail::uniform_distribution_double(
rocrand(state));
644__forceinline__ __device__ __host__
float rocrand_uniform(rocrand_state_lfsr113* state)
646 return rocrand_device::detail::uniform_distribution(
rocrand(state));
666 return rocrand_device::detail::uniform_distribution_double(
rocrand(state));
681__forceinline__ __device__ __host__
float rocrand_uniform(rocrand_state_threefry2x32_20* state)
683 return rocrand_device::detail::uniform_distribution(
rocrand(state));
701__forceinline__ __device__ __host__
double
704 return rocrand_device::detail::uniform_distribution_double(
rocrand(state));
719__forceinline__ __device__ __host__
float rocrand_uniform(rocrand_state_threefry2x64_20* state)
721 return rocrand_device::detail::uniform_distribution(
rocrand(state));
739__forceinline__ __device__ __host__
double
742 return rocrand_device::detail::uniform_distribution_double(
rocrand(state));
757__forceinline__ __device__ __host__
float rocrand_uniform(rocrand_state_threefry4x32_20* state)
759 return rocrand_device::detail::uniform_distribution(
rocrand(state));
777__forceinline__ __device__ __host__
double
780 return rocrand_device::detail::uniform_distribution_double(
rocrand(state));
795__forceinline__ __device__ __host__
float rocrand_uniform(rocrand_state_threefry4x64_20* state)
797 return rocrand_device::detail::uniform_distribution(
rocrand(state));
815__forceinline__ __device__ __host__
double
818 return rocrand_device::detail::uniform_distribution_double(
rocrand(state));
__forceinline__ __device__ __host__ float4 rocrand_uniform4(rocrand_state_philox4x32_10 *state)
Returns four uniformly distributed random float values from (0; 1] range.
Definition rocrand_uniform.h:276
__forceinline__ __device__ __host__ double4 rocrand_uniform_double4(rocrand_state_philox4x32_10 *state)
Returns four uniformly distributed random double values from (0; 1] range.
Definition rocrand_uniform.h:333
__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:386
__forceinline__ __device__ __host__ double rocrand_uniform_double(rocrand_state_philox4x32_10 *state)
Returns a uniformly distributed random double value from (0; 1] range.
Definition rocrand_uniform.h:294
__forceinline__ __device__ __host__ float rocrand_uniform(rocrand_state_philox4x32_10 *state)
Returns a uniformly distributed random float value from (0; 1] range.
Definition rocrand_uniform.h:236
__forceinline__ __device__ __host__ float2 rocrand_uniform2(rocrand_state_philox4x32_10 *state)
Returns two uniformly distributed random float values from (0; 1] range.
Definition rocrand_uniform.h:253
__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
__forceinline__ __device__ __host__ double2 rocrand_uniform_double2(rocrand_state_philox4x32_10 *state)
Returns two uniformly distributed random double values from (0; 1] range.
Definition rocrand_uniform.h:315