21#ifndef ROCRAND_DISCRETE_H_
22#define ROCRAND_DISCRETE_H_
26#include "rocrand/rocrand_lfsr113.h"
27#include "rocrand/rocrand_mrg31k3p.h"
28#include "rocrand/rocrand_mrg32k3a.h"
29#include "rocrand/rocrand_mtgp32.h"
30#include "rocrand/rocrand_philox4x32_10.h"
31#include "rocrand/rocrand_scrambled_sobol32.h"
32#include "rocrand/rocrand_scrambled_sobol64.h"
33#include "rocrand/rocrand_sobol32.h"
34#include "rocrand/rocrand_sobol64.h"
35#include "rocrand/rocrand_threefry2x32_20.h"
36#include "rocrand/rocrand_threefry2x64_20.h"
37#include "rocrand/rocrand_threefry4x32_20.h"
38#include "rocrand/rocrand_threefry4x64_20.h"
39#include "rocrand/rocrand_xorwow.h"
41#include "rocrand/rocrand_discrete_types.h"
46#if defined(__HIP_DEVICE_COMPILE__) && (defined(__GFX10__) || defined(__GFX11__))
47 #define ROCRAND_PREFER_CDF_OVER_ALIAS
58namespace rocrand_device {
61__forceinline__ __device__ __host__
unsigned int
62 discrete_alias(
const double x,
63 const unsigned int size,
64 const unsigned int offset,
65 const unsigned int* __restrict__ alias,
66 const double* __restrict__ probability)
71 const double nx = size * x;
72 const double fnx = floor(nx);
73 const double y = nx - fnx;
74 const unsigned int i =
static_cast<unsigned int>(fnx);
75 return offset + (y < probability[i] ? i : alias[i]);
78__forceinline__ __device__ __host__
unsigned int
84__forceinline__ __device__ __host__
unsigned int
87 constexpr double inv_double_32 = ROCRAND_2POW32_INV_DOUBLE;
88 const double x = r * inv_double_32;
89 return discrete_alias(x, dis);
93__forceinline__ __device__ __host__
unsigned int
96 constexpr double inv_double_32 = ROCRAND_2POW32_INV_DOUBLE;
97 const double x = r * inv_double_32;
98 return discrete_alias(x, dis);
101__forceinline__ __device__ __host__
unsigned int
104 constexpr double inv_double_64 = ROCRAND_2POW64_INV_DOUBLE;
105 const double x = r * inv_double_64;
106 return discrete_alias(x, dis);
109__forceinline__ __device__ __host__
unsigned int discrete_cdf(
const double x,
110 const unsigned int size,
111 const unsigned int offset,
112 const double* __restrict__ cdf)
116 unsigned int min = 0;
117 unsigned int max = size - 1;
120 const unsigned int center = (min + max) / 2;
121 const double p = cdf[center];
136__forceinline__ __device__ __host__
unsigned int
142__forceinline__ __device__ __host__
unsigned int
145 constexpr double inv_double_32 = ROCRAND_2POW32_INV_DOUBLE;
146 const double x = r * inv_double_32;
147 return discrete_cdf(x, dis);
151__forceinline__ __device__ __host__
unsigned int
154 constexpr double inv_double_32 = ROCRAND_2POW32_INV_DOUBLE;
155 const double x = r * inv_double_32;
156 return discrete_cdf(x, dis);
159__forceinline__ __device__ __host__
unsigned int
162 constexpr double inv_double_64 = ROCRAND_2POW64_INV_DOUBLE;
163 const double x = r * inv_double_64;
164 return discrete_cdf(x, dis);
187__forceinline__ __device__ __host__
unsigned int
191 return rocrand_device::detail::discrete_alias(
rocrand(state), *discrete_distribution);
211 rocrand_device::detail::discrete_alias(u4.x, *discrete_distribution),
212 rocrand_device::detail::discrete_alias(u4.y, *discrete_distribution),
213 rocrand_device::detail::discrete_alias(u4.z, *discrete_distribution),
214 rocrand_device::detail::discrete_alias(u4.w, *discrete_distribution)
230__forceinline__ __device__ __host__
unsigned int
234 return rocrand_device::detail::discrete_alias(
rocrand(state), *discrete_distribution);
249__forceinline__ __device__ __host__
unsigned int
253 return rocrand_device::detail::discrete_alias(
rocrand(state), *discrete_distribution);
268__forceinline__ __device__ __host__
unsigned int
272 return rocrand_device::detail::discrete_alias(
rocrand(state), *discrete_distribution);
287__forceinline__ __device__
unsigned int
291#ifdef ROCRAND_PREFER_CDF_OVER_ALIAS
292 return rocrand_device::detail::discrete_cdf(
rocrand(state), *discrete_distribution);
294 return rocrand_device::detail::discrete_alias(
rocrand(state), *discrete_distribution);
310__forceinline__ __device__ __host__
unsigned int
314 return rocrand_device::detail::discrete_cdf(
rocrand(state), *discrete_distribution);
329__forceinline__ __device__ __host__
unsigned int
333 return rocrand_device::detail::discrete_cdf(
rocrand(state), *discrete_distribution);
348__forceinline__ __device__ __host__
unsigned int
352 return rocrand_device::detail::discrete_cdf(
rocrand(state), *discrete_distribution);
367__forceinline__ __device__ __host__
unsigned int
371 return rocrand_device::detail::discrete_cdf(
rocrand(state), *discrete_distribution);
386__forceinline__ __device__ __host__
unsigned int
390#ifdef ROCRAND_PREFER_CDF_OVER_ALIAS
391 return rocrand_device::detail::discrete_cdf(
rocrand(state), *discrete_distribution);
393 return rocrand_device::detail::discrete_alias(
rocrand(state), *discrete_distribution);
409__forceinline__ __device__ __host__
unsigned int
413#ifdef ROCRAND_PREFER_CDF_OVER_ALIAS
414 return rocrand_device::detail::discrete_cdf(
rocrand(state), *discrete_distribution);
416 return rocrand_device::detail::discrete_alias(
rocrand(state), *discrete_distribution);
432__forceinline__ __device__ __host__
unsigned int
436#ifdef ROCRAND_PREFER_CDF_OVER_ALIAS
437 return rocrand_device::detail::discrete_cdf(
rocrand(state), *discrete_distribution);
439 return rocrand_device::detail::discrete_alias(
rocrand(state), *discrete_distribution);
455__forceinline__ __device__ __host__
unsigned int
459#ifdef ROCRAND_PREFER_CDF_OVER_ALIAS
460 return rocrand_device::detail::discrete_cdf(
rocrand(state), *discrete_distribution);
462 return rocrand_device::detail::discrete_alias(
rocrand(state), *discrete_distribution);
478__forceinline__ __device__ __host__
unsigned int
482#ifdef ROCRAND_PREFER_CDF_OVER_ALIAS
483 return rocrand_device::detail::discrete_cdf(
rocrand(state), *discrete_distribution);
485 return rocrand_device::detail::discrete_alias(
rocrand(state), *discrete_distribution);
492#if defined(ROCRAND_PREFER_CDF_OVER_ALIAS)
493 #undef ROCRAND_PREFER_CDF_OVER_ALIAS
__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__ uint4 rocrand_discrete4(rocrand_state_philox4x32_10 *state, const rocrand_discrete_distribution discrete_distribution)
Returns four discrete distributed unsigned int values.
Definition rocrand_discrete.h:206
__forceinline__ __device__ __host__ unsigned int rocrand_discrete(rocrand_state_philox4x32_10 *state, const rocrand_discrete_distribution discrete_distribution)
Returns a discrete distributed unsigned int value.
Definition rocrand_discrete.h:188
__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
Represents a discrete probability distribution.
Definition rocrand_discrete_types.h:26
unsigned int size
Number of entries in the probability table.
Definition rocrand_discrete_types.h:28
unsigned int * alias
Alias table.
Definition rocrand_discrete_types.h:33
double * cdf
Cumulative distribution function.
Definition rocrand_discrete_types.h:38
double * probability
Probability data for the alias table.
Definition rocrand_discrete_types.h:35
unsigned int offset
The distribution can be offset.
Definition rocrand_discrete_types.h:30