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

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

API library: /build/reproducible-path/rocrand-6.4.1/library/include/rocrand/rocrand_discrete.h Source File
rocrand_discrete.h
1// Copyright (c) 2017-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_DISCRETE_H_
22#define ROCRAND_DISCRETE_H_
23
24#include <math.h>
25
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"
40
41#include "rocrand/rocrand_discrete_types.h"
42
43// On certain architectures such as NAVI2 and NAVI3, double arithmetic is significantly slower.
44// In such cases we want to prefer the CDF method over the alias method.
45// This macro is undefined at the end of the file.
46#if defined(__HIP_DEVICE_COMPILE__) && (defined(__GFX10__) || defined(__GFX11__))
47 #define ROCRAND_PREFER_CDF_OVER_ALIAS
48#endif
49
50// Alias method
51//
52// Walker, A. J.
53// An Efficient Method for Generating Discrete Random Variables with General Distributions, 1977
54//
55// Vose M. D.
56// A Linear Algorithm For Generating Random Numbers With a Given Distribution, 1991
57
58namespace rocrand_device {
59namespace detail {
60
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)
67{
68 // Calculate value using Alias table
69
70 // x is [0, 1)
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]);
76}
77
78__forceinline__ __device__ __host__ unsigned int
79 discrete_alias(const double x, const rocrand_discrete_distribution_st& dis)
80{
81 return discrete_alias(x, dis.size, dis.offset, dis.alias, dis.probability);
82}
83
84__forceinline__ __device__ __host__ unsigned int
85 discrete_alias(const unsigned int r, const rocrand_discrete_distribution_st& dis)
86{
87 constexpr double inv_double_32 = ROCRAND_2POW32_INV_DOUBLE;
88 const double x = r * inv_double_32;
89 return discrete_alias(x, dis);
90}
91
92// To prevent ambiguity compile error when compiler is facing the type "unsigned long"!!!
93__forceinline__ __device__ __host__ unsigned int
94 discrete_alias(const unsigned long r, const rocrand_discrete_distribution_st& dis)
95{
96 constexpr double inv_double_32 = ROCRAND_2POW32_INV_DOUBLE;
97 const double x = r * inv_double_32;
98 return discrete_alias(x, dis);
99}
100
101__forceinline__ __device__ __host__ unsigned int
102 discrete_alias(const unsigned long long int r, const rocrand_discrete_distribution_st& dis)
103{
104 constexpr double inv_double_64 = ROCRAND_2POW64_INV_DOUBLE;
105 const double x = r * inv_double_64;
106 return discrete_alias(x, dis);
107}
108
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)
113{
114 // Calculate value using binary search in CDF
115
116 unsigned int min = 0;
117 unsigned int max = size - 1;
118 do
119 {
120 const unsigned int center = (min + max) / 2;
121 const double p = cdf[center];
122 if(x > p)
123 {
124 min = center + 1;
125 }
126 else
127 {
128 max = center;
129 }
130 }
131 while(min != max);
132
133 return offset + min;
134}
135
136__forceinline__ __device__ __host__ unsigned int
137 discrete_cdf(const double x, const rocrand_discrete_distribution_st& dis)
138{
139 return discrete_cdf(x, dis.size, dis.offset, dis.cdf);
140}
141
142__forceinline__ __device__ __host__ unsigned int
143 discrete_cdf(const unsigned int r, const rocrand_discrete_distribution_st& dis)
144{
145 constexpr double inv_double_32 = ROCRAND_2POW32_INV_DOUBLE;
146 const double x = r * inv_double_32;
147 return discrete_cdf(x, dis);
148}
149
150// To prevent ambiguity compile error when compiler is facing the type "unsigned long"!!!
151__forceinline__ __device__ __host__ unsigned int
152 discrete_cdf(const unsigned long r, const rocrand_discrete_distribution_st& dis)
153{
154 constexpr double inv_double_32 = ROCRAND_2POW32_INV_DOUBLE;
155 const double x = r * inv_double_32;
156 return discrete_cdf(x, dis);
157}
158
159__forceinline__ __device__ __host__ unsigned int
160 discrete_cdf(const unsigned long long int r, const rocrand_discrete_distribution_st& dis)
161{
162 constexpr double inv_double_64 = ROCRAND_2POW64_INV_DOUBLE;
163 const double x = r * inv_double_64;
164 return discrete_cdf(x, dis);
165}
166
167} // end namespace detail
168} // end namespace rocrand_device
169
187__forceinline__ __device__ __host__ unsigned int
188 rocrand_discrete(rocrand_state_philox4x32_10* state,
189 const rocrand_discrete_distribution discrete_distribution)
190{
191 return rocrand_device::detail::discrete_alias(rocrand(state), *discrete_distribution);
192}
193
206__forceinline__ __device__ __host__ uint4 rocrand_discrete4(
207 rocrand_state_philox4x32_10* state, const rocrand_discrete_distribution discrete_distribution)
208{
209 const uint4 u4 = rocrand4(state);
210 return uint4 {
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)
215 };
216}
217
230__forceinline__ __device__ __host__ unsigned int
231 rocrand_discrete(rocrand_state_mrg31k3p* state,
232 const rocrand_discrete_distribution discrete_distribution)
233{
234 return rocrand_device::detail::discrete_alias(rocrand(state), *discrete_distribution);
235}
236
249__forceinline__ __device__ __host__ unsigned int
250 rocrand_discrete(rocrand_state_mrg32k3a* state,
251 const rocrand_discrete_distribution discrete_distribution)
252{
253 return rocrand_device::detail::discrete_alias(rocrand(state), *discrete_distribution);
254}
255
268__forceinline__ __device__ __host__ unsigned int
269 rocrand_discrete(rocrand_state_xorwow* state,
270 const rocrand_discrete_distribution discrete_distribution)
271{
272 return rocrand_device::detail::discrete_alias(rocrand(state), *discrete_distribution);
273}
274
287__forceinline__ __device__ unsigned int
288 rocrand_discrete(rocrand_state_mtgp32* state,
289 const rocrand_discrete_distribution discrete_distribution)
290{
291#ifdef ROCRAND_PREFER_CDF_OVER_ALIAS
292 return rocrand_device::detail::discrete_cdf(rocrand(state), *discrete_distribution);
293#else
294 return rocrand_device::detail::discrete_alias(rocrand(state), *discrete_distribution);
295#endif
296}
297
310__forceinline__ __device__ __host__ unsigned int
311 rocrand_discrete(rocrand_state_sobol32* state,
312 const rocrand_discrete_distribution discrete_distribution)
313{
314 return rocrand_device::detail::discrete_cdf(rocrand(state), *discrete_distribution);
315}
316
329__forceinline__ __device__ __host__ unsigned int
330 rocrand_discrete(rocrand_state_scrambled_sobol32* state,
331 const rocrand_discrete_distribution discrete_distribution)
332{
333 return rocrand_device::detail::discrete_cdf(rocrand(state), *discrete_distribution);
334}
335
348__forceinline__ __device__ __host__ unsigned int
349 rocrand_discrete(rocrand_state_sobol64* state,
350 const rocrand_discrete_distribution discrete_distribution)
351{
352 return rocrand_device::detail::discrete_cdf(rocrand(state), *discrete_distribution);
353}
354
367__forceinline__ __device__ __host__ unsigned int
368 rocrand_discrete(rocrand_state_scrambled_sobol64* state,
369 const rocrand_discrete_distribution discrete_distribution)
370{
371 return rocrand_device::detail::discrete_cdf(rocrand(state), *discrete_distribution);
372}
373
386__forceinline__ __device__ __host__ unsigned int
387 rocrand_discrete(rocrand_state_lfsr113* state,
388 const rocrand_discrete_distribution discrete_distribution)
389{
390#ifdef ROCRAND_PREFER_CDF_OVER_ALIAS
391 return rocrand_device::detail::discrete_cdf(rocrand(state), *discrete_distribution);
392#else
393 return rocrand_device::detail::discrete_alias(rocrand(state), *discrete_distribution);
394#endif
395}
396
409__forceinline__ __device__ __host__ unsigned int
410 rocrand_discrete(rocrand_state_threefry2x32_20* state,
411 const rocrand_discrete_distribution discrete_distribution)
412{
413#ifdef ROCRAND_PREFER_CDF_OVER_ALIAS
414 return rocrand_device::detail::discrete_cdf(rocrand(state), *discrete_distribution);
415#else
416 return rocrand_device::detail::discrete_alias(rocrand(state), *discrete_distribution);
417#endif
418}
419
432__forceinline__ __device__ __host__ unsigned int
433 rocrand_discrete(rocrand_state_threefry2x64_20* state,
434 const rocrand_discrete_distribution discrete_distribution)
435{
436#ifdef ROCRAND_PREFER_CDF_OVER_ALIAS
437 return rocrand_device::detail::discrete_cdf(rocrand(state), *discrete_distribution);
438#else
439 return rocrand_device::detail::discrete_alias(rocrand(state), *discrete_distribution);
440#endif
441}
442
455__forceinline__ __device__ __host__ unsigned int
456 rocrand_discrete(rocrand_state_threefry4x32_20* state,
457 const rocrand_discrete_distribution discrete_distribution)
458{
459#ifdef ROCRAND_PREFER_CDF_OVER_ALIAS
460 return rocrand_device::detail::discrete_cdf(rocrand(state), *discrete_distribution);
461#else
462 return rocrand_device::detail::discrete_alias(rocrand(state), *discrete_distribution);
463#endif
464}
465
478__forceinline__ __device__ __host__ unsigned int
479 rocrand_discrete(rocrand_state_threefry4x64_20* state,
480 const rocrand_discrete_distribution discrete_distribution)
481{
482#ifdef ROCRAND_PREFER_CDF_OVER_ALIAS
483 return rocrand_device::detail::discrete_cdf(rocrand(state), *discrete_distribution);
484#else
485 return rocrand_device::detail::discrete_alias(rocrand(state), *discrete_distribution);
486#endif
487}
488
// end of group rocranddevice
490
491// Undefine the macro that may be defined at the top of the file!
492#if defined(ROCRAND_PREFER_CDF_OVER_ALIAS)
493 #undef ROCRAND_PREFER_CDF_OVER_ALIAS
494#endif
495
496#endif // ROCRAND_DISCRETE_H_
__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