Skip to main content
`Ctrl`+`K`

# /build/rocrand-cWAlCE/rocrand-7.2.4/library/include/rocrand/rocrand_discrete.h Source File

API library: /build/rocrand-cWAlCE/rocrand-7.2.4/library/include/rocrand/rocrand_discrete.h Source File
rocrand_discrete.h
1// Copyright (c) 2017-2025 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 "rocrand/rocrand_common.h"
25#include "rocrand/rocrand_lfsr113.h"
26#include "rocrand/rocrand_mrg31k3p.h"
27#include "rocrand/rocrand_mrg32k3a.h"
28#include "rocrand/rocrand_mtgp32.h"
29#include "rocrand/rocrand_philox4x32_10.h"
30#include "rocrand/rocrand_scrambled_sobol32.h"
31#include "rocrand/rocrand_scrambled_sobol64.h"
32#include "rocrand/rocrand_sobol32.h"
33#include "rocrand/rocrand_sobol64.h"
34#include "rocrand/rocrand_threefry2x32_20.h"
35#include "rocrand/rocrand_threefry2x64_20.h"
36#include "rocrand/rocrand_threefry4x32_20.h"
37#include "rocrand/rocrand_threefry4x64_20.h"
38#include "rocrand/rocrand_xorwow.h"
39
40#include <hip/hip_runtime.h>
41
42#include <math.h>
43
44// On certain architectures such as NAVI2 and NAVI3, double arithmetic is significantly slower.
45// In such cases we want to prefer the CDF method over the alias method.
46// This macro is undefined at the end of the file.
47#if defined(__HIP_DEVICE_COMPILE__) && (defined(__GFX10__) || defined(__GFX11__))
48 #define ROCRAND_PREFER_CDF_OVER_ALIAS
49#endif
50
51// Alias method
52//
53// Walker, A. J.
54// An Efficient Method for Generating Discrete Random Variables with General Distributions, 1977
55//
56// Vose M. D.
57// A Linear Algorithm For Generating Random Numbers With a Given Distribution, 1991
58
59 namespace rocrand_device {
60 namespace detail {
61
62__forceinline__ __device__ __host__ unsigned int
63 discrete_alias(const double x,
64 const unsigned int size,
65 const unsigned int offset,
66 const unsigned int* __restrict__ alias,
67 const double* __restrict__ probability)
68{
69 // Calculate value using Alias table
70
71 // x is [0, 1)
72 const double nx = size * x;
73 const double fnx = floor(nx);
74 const double y = nx - fnx;
75 const unsigned int i = static_cast<unsigned int>(fnx);
76 return offset + (y < probability[i] ? i : alias[i]);
77}
78
79__forceinline__ __device__ __host__ unsigned int
80 discrete_alias(const double x, const rocrand_discrete_distribution_st& dis)
81{
82 return discrete_alias(x, dis.size, dis.offset, dis.alias, dis.probability);
83}
84
85__forceinline__ __device__ __host__ unsigned int
86 discrete_alias(const unsigned int r, const rocrand_discrete_distribution_st& dis)
87{
88 constexpr double inv_double_32 = ROCRAND_2POW32_INV_DOUBLE;
89 const double x = r * inv_double_32;
90 return discrete_alias(x, dis);
91}
92
93// To prevent ambiguity compile error when compiler is facing the type "unsigned long"!!!
94__forceinline__ __device__ __host__ unsigned int
95 discrete_alias(const unsigned long r, const rocrand_discrete_distribution_st& dis)
96{
97 constexpr double inv_double_32 = ROCRAND_2POW32_INV_DOUBLE;
98 const double x = r * inv_double_32;
99 return discrete_alias(x, dis);
100}
101
102__forceinline__ __device__ __host__ unsigned int
103 discrete_alias(const unsigned long long int r, const rocrand_discrete_distribution_st& dis)
104{
105 constexpr double inv_double_64 = ROCRAND_2POW64_INV_DOUBLE;
106 const double x = r * inv_double_64;
107 return discrete_alias(x, dis);
108}
109
110__forceinline__ __device__ __host__ unsigned int discrete_cdf(const double x,
111 const unsigned int size,
112 const unsigned int offset,
113 const double* __restrict__ cdf)
114{
115 // Calculate value using binary search in CDF
116
117 unsigned int min = 0;
118 unsigned int max = size - 1;
119 do
120 {
121 const unsigned int center = (min + max) / 2;
122 const double p = cdf[center];
123 if(x > p)
124 {
125 min = center + 1;
126 }
127 else
128 {
129 max = center;
130 }
131 }
132 while(min != max);
133
134 return offset + min;
135}
136
137__forceinline__ __device__ __host__ unsigned int
138 discrete_cdf(const double x, const rocrand_discrete_distribution_st& dis)
139{
140 return discrete_cdf(x, dis.size, dis.offset, dis.cdf);
141}
142
143__forceinline__ __device__ __host__ unsigned int
144 discrete_cdf(const unsigned int r, const rocrand_discrete_distribution_st& dis)
145{
146 constexpr double inv_double_32 = ROCRAND_2POW32_INV_DOUBLE;
147 const double x = r * inv_double_32;
148 return discrete_cdf(x, dis);
149}
150
151// To prevent ambiguity compile error when compiler is facing the type "unsigned long"!!!
152__forceinline__ __device__ __host__ unsigned int
153 discrete_cdf(const unsigned long r, const rocrand_discrete_distribution_st& dis)
154{
155 constexpr double inv_double_32 = ROCRAND_2POW32_INV_DOUBLE;
156 const double x = r * inv_double_32;
157 return discrete_cdf(x, dis);
158}
159
160__forceinline__ __device__ __host__ unsigned int
161 discrete_cdf(const unsigned long long int r, const rocrand_discrete_distribution_st& dis)
162{
163 constexpr double inv_double_64 = ROCRAND_2POW64_INV_DOUBLE;
164 const double x = r * inv_double_64;
165 return discrete_cdf(x, dis);
166}
167
168} // end namespace detail
169} // end namespace rocrand_device
170
175
188__forceinline__ __device__ __host__
189 unsigned int rocrand_discrete(rocrand_state_philox4x32_10* state,
190 const rocrand_discrete_distribution discrete_distribution)
191{
192 return rocrand_device::detail::discrete_alias(rocrand(state), *discrete_distribution);
193}
194
207__forceinline__ __device__ __host__
208uint4 rocrand_discrete4(rocrand_state_philox4x32_10* state,
209 const rocrand_discrete_distribution discrete_distribution)
210{
211 const uint4 u4 = rocrand4(state);
212 return uint4 {
213 rocrand_device::detail::discrete_alias(u4.x, *discrete_distribution),
214 rocrand_device::detail::discrete_alias(u4.y, *discrete_distribution),
215 rocrand_device::detail::discrete_alias(u4.z, *discrete_distribution),
216 rocrand_device::detail::discrete_alias(u4.w, *discrete_distribution)
217 };
218}
219
232__forceinline__ __device__ __host__
233 unsigned int rocrand_discrete(rocrand_state_mrg31k3p* state,
234 const rocrand_discrete_distribution discrete_distribution)
235{
236 return rocrand_device::detail::discrete_alias(rocrand(state), *discrete_distribution);
237}
238
251__forceinline__ __device__ __host__
252 unsigned int rocrand_discrete(rocrand_state_mrg32k3a* state,
253 const rocrand_discrete_distribution discrete_distribution)
254{
255 return rocrand_device::detail::discrete_alias(rocrand(state), *discrete_distribution);
256}
257
270__forceinline__ __device__ __host__
271 unsigned int rocrand_discrete(rocrand_state_xorwow* state,
272 const rocrand_discrete_distribution discrete_distribution)
273{
274 return rocrand_device::detail::discrete_alias(rocrand(state), *discrete_distribution);
275}
276
289__forceinline__ __device__
290 unsigned int rocrand_discrete(rocrand_state_mtgp32* state,
291 const rocrand_discrete_distribution discrete_distribution)
292{
293#ifdef ROCRAND_PREFER_CDF_OVER_ALIAS
294 return rocrand_device::detail::discrete_cdf(rocrand(state), *discrete_distribution);
295#else
296 return rocrand_device::detail::discrete_alias(rocrand(state), *discrete_distribution);
297#endif
298}
299
312__forceinline__ __device__ __host__
313 unsigned int rocrand_discrete(rocrand_state_sobol32* state,
314 const rocrand_discrete_distribution discrete_distribution)
315{
316 return rocrand_device::detail::discrete_cdf(rocrand(state), *discrete_distribution);
317}
318
331__forceinline__ __device__ __host__
332 unsigned int rocrand_discrete(rocrand_state_scrambled_sobol32* state,
333 const rocrand_discrete_distribution discrete_distribution)
334{
335 return rocrand_device::detail::discrete_cdf(rocrand(state), *discrete_distribution);
336}
337
350__forceinline__ __device__ __host__
351 unsigned int rocrand_discrete(rocrand_state_sobol64* state,
352 const rocrand_discrete_distribution discrete_distribution)
353{
354 return rocrand_device::detail::discrete_cdf(rocrand(state), *discrete_distribution);
355}
356
369__forceinline__ __device__ __host__
370 unsigned int rocrand_discrete(rocrand_state_scrambled_sobol64* state,
371 const rocrand_discrete_distribution discrete_distribution)
372{
373 return rocrand_device::detail::discrete_cdf(rocrand(state), *discrete_distribution);
374}
375
388__forceinline__ __device__ __host__
389 unsigned int rocrand_discrete(rocrand_state_lfsr113* state,
390 const rocrand_discrete_distribution discrete_distribution)
391{
392#ifdef ROCRAND_PREFER_CDF_OVER_ALIAS
393 return rocrand_device::detail::discrete_cdf(rocrand(state), *discrete_distribution);
394#else
395 return rocrand_device::detail::discrete_alias(rocrand(state), *discrete_distribution);
396#endif
397}
398
411__forceinline__ __device__ __host__
412 unsigned int rocrand_discrete(rocrand_state_threefry2x32_20* state,
413 const rocrand_discrete_distribution discrete_distribution)
414{
415#ifdef ROCRAND_PREFER_CDF_OVER_ALIAS
416 return rocrand_device::detail::discrete_cdf(rocrand(state), *discrete_distribution);
417#else
418 return rocrand_device::detail::discrete_alias(rocrand(state), *discrete_distribution);
419#endif
420}
421
434__forceinline__ __device__ __host__
435 unsigned int rocrand_discrete(rocrand_state_threefry2x64_20* state,
436 const rocrand_discrete_distribution discrete_distribution)
437{
438#ifdef ROCRAND_PREFER_CDF_OVER_ALIAS
439 return rocrand_device::detail::discrete_cdf(rocrand(state), *discrete_distribution);
440#else
441 return rocrand_device::detail::discrete_alias(rocrand(state), *discrete_distribution);
442#endif
443}
444
457__forceinline__ __device__ __host__
458 unsigned int rocrand_discrete(rocrand_state_threefry4x32_20* state,
459 const rocrand_discrete_distribution discrete_distribution)
460{
461#ifdef ROCRAND_PREFER_CDF_OVER_ALIAS
462 return rocrand_device::detail::discrete_cdf(rocrand(state), *discrete_distribution);
463#else
464 return rocrand_device::detail::discrete_alias(rocrand(state), *discrete_distribution);
465#endif
466}
467
480__forceinline__ __device__ __host__
481 unsigned int rocrand_discrete(rocrand_state_threefry4x64_20* state,
482 const rocrand_discrete_distribution discrete_distribution)
483{
484#ifdef ROCRAND_PREFER_CDF_OVER_ALIAS
485 return rocrand_device::detail::discrete_cdf(rocrand(state), *discrete_distribution);
486#else
487 return rocrand_device::detail::discrete_alias(rocrand(state), *discrete_distribution);
488#endif
489}
490 // end of group rocranddevice
492
493// Undefine the macro that may be defined at the top of the file!
494#if defined(ROCRAND_PREFER_CDF_OVER_ALIAS)
495 #undef ROCRAND_PREFER_CDF_OVER_ALIAS
496#endif
497
498#endif // ROCRAND_DISCRETE_H_
rocrand4
__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:377
rocrand_discrete4
__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:208
rocrand_discrete
__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:189
rocrand
__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:277
rocrand_discrete_distribution_st::size
unsigned int size
Number of entries in the probability table.
Definition rocrand_discrete_types.h:28
rocrand_discrete_distribution_st::alias
unsigned int * alias
Alias table.
Definition rocrand_discrete_types.h:33
rocrand_discrete_distribution_st::cdf
double * cdf
Cumulative distribution function.
Definition rocrand_discrete_types.h:38
rocrand_discrete_distribution_st::probability
double * probability
Probability data for the alias table.
Definition rocrand_discrete_types.h:35
rocrand_discrete_distribution_st::offset
unsigned int offset
The distribution can be offset.
Definition rocrand_discrete_types.h:30