/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-rocrand/checkouts/latest/library/include/rocrand/rocrand_discrete.h Source File

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-rocrand/checkouts/latest/library/include/rocrand/rocrand_discrete.h Source File#

API library: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-rocrand/checkouts/latest/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 
58 namespace rocrand_device {
59 namespace 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