/build/rocrand-A0YaxB/rocrand-7.14/library/include/rocrand/rocrand_poisson.h Source File

/build/rocrand-A0YaxB/rocrand-7.14/library/include/rocrand/rocrand_poisson.h Source File#

API library: /build/rocrand-A0YaxB/rocrand-7.14/library/include/rocrand/rocrand_poisson.h Source File
rocrand_poisson.h
1// Copyright (c) 2017-2026 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_POISSON_H_
22#define ROCRAND_POISSON_H_
23
28
29#include <math.h>
30
31#include "rocrand/rocrand_lfsr113.h"
32#include "rocrand/rocrand_mrg31k3p.h"
33#include "rocrand/rocrand_mrg32k3a.h"
34#include "rocrand/rocrand_mtgp32.h"
35#include "rocrand/rocrand_philox4x32_10.h"
36#include "rocrand/rocrand_scrambled_sobol32.h"
37#include "rocrand/rocrand_scrambled_sobol64.h"
38#include "rocrand/rocrand_sobol32.h"
39#include "rocrand/rocrand_sobol64.h"
40#include "rocrand/rocrand_threefry2x32_20.h"
41#include "rocrand/rocrand_threefry2x64_20.h"
42#include "rocrand/rocrand_threefry4x32_20.h"
43#include "rocrand/rocrand_threefry4x64_20.h"
44#include "rocrand/rocrand_xorwow.h"
45
46#include "rocrand/rocrand_normal.h"
47#include "rocrand/rocrand_uniform.h"
48
49#include <hip/hip_runtime.h>
50
51namespace rocrand_device
52{
53namespace detail
54{
55
56constexpr double lambda_threshold_small = 64.0;
57constexpr double lambda_threshold_huge = 4000.0;
58
59template<class State, typename Result_Type = unsigned int>
60__forceinline__ __device__ __host__
61Result_Type poisson_distribution_small(State& state, double lambda)
62
63{
64 // Knuth's method
65
66 const double limit = exp(-lambda);
67 Result_Type k = 0;
68 double product = 1.0;
69
70 do
71 {
72 k++;
73 product *= rocrand_uniform_double(state);
74 }
75 while(product > limit);
76
77 return k - 1;
78}
79
80__forceinline__ __device__ __host__
81double lgamma_approx(const double x)
82{
83 // Lanczos approximation (g = 7, n = 9)
84
85 const double z = x - 1.0;
86
87 const int g = 7;
88 const int n = 9;
89 const double coefs[n] = {0.99999999999980993227684700473478,
90 676.520368121885098567009190444019,
91 -1259.13921672240287047156078755283,
92 771.3234287776530788486528258894,
93 -176.61502916214059906584551354,
94 12.507343278686904814458936853,
95 -0.13857109526572011689554707,
96 9.984369578019570859563e-6,
97 1.50563273514931155834e-7};
98 double sum = 0.0;
99#pragma unroll
100 for(int i = n - 1; i > 0; i--)
101
102 {
103 sum += coefs[i] / (z + i);
104 }
105 sum += coefs[0];
106
107 const double log_sqrt_2_pi = 0.9189385332046727418;
108 const double e = 2.718281828459045090796;
109 return (log_sqrt_2_pi + log(sum) - g) + (z + 0.5) * log((z + g + 0.5) / e);
110}
111
112template<class State, typename Result_Type = unsigned int>
113__forceinline__ __device__ __host__
114Result_Type poisson_distribution_large(State& state, double lambda)
115{
116 // Rejection method PA, A. C. Atkinson
117
118 const double c = 0.767 - 3.36 / lambda;
119 const double beta = ROCRAND_PI_DOUBLE / sqrt(3.0 * lambda);
120 const double alpha = beta * lambda;
121 const double k = log(c) - lambda - log(beta);
122 const double log_lambda = log(lambda);
123
124 while(true)
125 {
126 const double u = rocrand_uniform_double(state);
127 const double x = (alpha - log((1.0 - u) / u)) / beta;
128 const double n = floor(x + 0.5);
129 if(n < 0)
130 {
131 continue;
132 }
133 const double v = rocrand_uniform_double(state);
134 const double y = alpha - beta * x;
135 const double t = 1.0 + exp(y);
136 const double lhs = y + log(v / (t * t));
137 const double rhs = k + n * log_lambda - lgamma_approx(n + 1.0);
138 if(lhs <= rhs)
139 {
140 return static_cast<Result_Type>(n);
141 }
142 }
143}
144
145template<class State, typename Result_Type = unsigned int>
146__forceinline__ __device__ __host__
147Result_Type poisson_distribution_huge(State& state, double lambda)
148{
149 // Approximate Poisson distribution with normal distribution
150
151 const double n = rocrand_normal_double(state);
152 return static_cast<Result_Type>(round(sqrt(lambda) * n + lambda));
153}
154
155template<class State, typename Result_Type = unsigned int>
156__forceinline__ __device__ __host__
157Result_Type poisson_distribution(State& state, double lambda)
158{
159 if(lambda < lambda_threshold_small)
160 {
161 return poisson_distribution_small<State, Result_Type>(state, lambda);
162 }
163 else if(lambda <= lambda_threshold_huge)
164 {
165 return poisson_distribution_large<State, Result_Type>(state, lambda);
166 }
167 else
168 {
169 return poisson_distribution_huge<State, Result_Type>(state, lambda);
170 }
171}
172
173template<class State, typename Result_Type = unsigned int>
174__forceinline__ __device__ __host__
175Result_Type poisson_distribution_itr(State& state, double lambda)
176{
177 // Algorithm ITR
178 // George S. Fishman
179 // Discrete-Event Simulation: Modeling, Programming, and Analysis
180 // p. 333
181 Result_Type k = 0;
182 // Algorithm ITR uses u from (0, 1) and uniform_double returns (0, 1]
183 // Change u to ensure that 1 is never generated,
184 // otherwise the inner loop never ends.
185 const double u = rocrand_uniform_double(state) - ROCRAND_2POW32_INV_DOUBLE / 2.0;
186 const double L = (lambda > 500.0) ? exp(-500.0) : exp(-lambda);
187
188 // First iteration, pow = 0 < lambda
189 double x = L;
190 double y = L;
191 while(u > y)
192 {
193 k++;
194 x *= ((double)lambda / (double)k);
195 y += x;
196 }
197
198 // Second iteration only if lambda > 500.0. No more iterations needed
199 // as we only call this method with lambda < 1000.0
200 if((double)500 < lambda)
201 {
202 // Now we now that necessarily lambda > upow = 500.0, so we already calculated before L = exp(-500)
203 x *= L;
204 y *= L;
205 while(u > y)
206 {
207 k++;
208 x *= ((double)lambda / (double)k);
209 y += x;
210 }
211 }
212 return k;
213}
214
215template<class State, typename Result_Type = unsigned int>
216__forceinline__ __device__ __host__
217Result_Type poisson_distribution_inv(State& state, double lambda)
218{
219 if(lambda < 1000.0)
220 {
221 return poisson_distribution_itr<State, Result_Type>(state, lambda);
222 }
223 else
224 {
225 return poisson_distribution_huge<State, Result_Type>(state, lambda);
226 }
227}
228
229} // end namespace detail
230} // end namespace rocrand_device
231
243#ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE
244__forceinline__ __device__ __host__
245unsigned int rocrand_poisson(rocrand_state_philox4x32_10* state, double lambda)
246{
247 return rocrand_device::detail::poisson_distribution<rocrand_state_philox4x32_10*, unsigned int>(
248 state,
249 lambda);
250}
251
263__forceinline__ __device__ __host__
264uint4 rocrand_poisson4(rocrand_state_philox4x32_10* state, double lambda)
265{
266 return uint4{
267 rocrand_device::detail::poisson_distribution<rocrand_state_philox4x32_10*, unsigned int>(
268 state,
269 lambda),
270 rocrand_device::detail::poisson_distribution<rocrand_state_philox4x32_10*, unsigned int>(
271 state,
272 lambda),
273 rocrand_device::detail::poisson_distribution<rocrand_state_philox4x32_10*, unsigned int>(
274 state,
275 lambda),
276 rocrand_device::detail::poisson_distribution<rocrand_state_philox4x32_10*, unsigned int>(
277 state,
278 lambda)};
279}
280#endif // ROCRAND_DETAIL_BM_NOT_IN_STATE
281
293#ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE
294__forceinline__ __device__ __host__
295unsigned int rocrand_poisson(rocrand_state_mrg31k3p* state, double lambda)
296{
297 return rocrand_device::detail::poisson_distribution<rocrand_state_mrg31k3p*, unsigned int>(
298 state,
299 lambda);
300}
301#endif // ROCRAND_DETAIL_BM_NOT_IN_STATE
302
314#ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE
315__forceinline__ __device__ __host__
316unsigned int rocrand_poisson(rocrand_state_mrg32k3a* state, double lambda)
317{
318 return rocrand_device::detail::poisson_distribution<rocrand_state_mrg32k3a*, unsigned int>(
319 state,
320 lambda);
321}
322#endif // ROCRAND_DETAIL_BM_NOT_IN_STATE
323
335#ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE
336__forceinline__ __device__ __host__
337unsigned int rocrand_poisson(rocrand_state_xorwow* state, double lambda)
338{
339 return rocrand_device::detail::poisson_distribution<rocrand_state_xorwow*, unsigned int>(
340 state,
341 lambda);
342}
343#endif // ROCRAND_DETAIL_BM_NOT_IN_STATE
344
356__forceinline__ __device__ __host__
357unsigned int rocrand_poisson(rocrand_state_mtgp32* state, double lambda)
358{
359 return rocrand_device::detail::poisson_distribution_inv<rocrand_state_mtgp32*, unsigned int>(
360 state,
361 lambda);
362}
363
375__forceinline__ __device__ __host__
376unsigned int rocrand_poisson(rocrand_state_sobol32* state, double lambda)
377{
378 return rocrand_device::detail::poisson_distribution_inv<rocrand_state_sobol32*, unsigned int>(
379 state,
380 lambda);
381}
382
394__forceinline__ __device__ __host__
395unsigned int rocrand_poisson(rocrand_state_scrambled_sobol32* state, double lambda)
396{
397 return rocrand_device::detail::poisson_distribution_inv<rocrand_state_scrambled_sobol32*,
398 unsigned int>(state, lambda);
399}
400
412__forceinline__ __device__ __host__
413unsigned int rocrand_poisson(rocrand_state_sobol64* state, double lambda)
414{
415 return rocrand_device::detail::poisson_distribution_inv<rocrand_state_sobol64*, unsigned int>(
416 state,
417 lambda);
418}
419
431__forceinline__ __device__ __host__
432unsigned int rocrand_poisson(rocrand_state_scrambled_sobol64* state, double lambda)
433{
434 return rocrand_device::detail::poisson_distribution_inv<rocrand_state_scrambled_sobol64*,
435 unsigned int>(state, lambda);
436}
437
449__forceinline__ __device__ __host__
450unsigned int rocrand_poisson(rocrand_state_lfsr113* state, double lambda)
451{
452 return rocrand_device::detail::poisson_distribution_inv<rocrand_state_lfsr113*, unsigned int>(
453 state,
454 lambda);
455}
456
468__forceinline__ __device__ __host__
469unsigned int rocrand_poisson(rocrand_state_threefry2x32_20* state, double lambda)
470{
471 return rocrand_device::detail::poisson_distribution_inv(state, lambda);
472}
473
485__forceinline__ __device__ __host__
486unsigned int rocrand_poisson(rocrand_state_threefry2x64_20* state, double lambda)
487{
488 return rocrand_device::detail::poisson_distribution_inv(state, lambda);
489}
490
502__forceinline__ __device__ __host__
503unsigned int rocrand_poisson(rocrand_state_threefry4x32_20* state, double lambda)
504{
505 return rocrand_device::detail::poisson_distribution_inv(state, lambda);
506}
507
519__forceinline__ __device__ __host__
520unsigned int rocrand_poisson(rocrand_state_threefry4x64_20* state, double lambda)
521{
522 return rocrand_device::detail::poisson_distribution_inv(state, lambda);
523}
524 // end of group rocranddevice
526
527#endif // ROCRAND_POISSON_H_
__forceinline__ __device__ __host__ unsigned int rocrand_poisson(rocrand_state_philox4x32_10 *state, double lambda)
Returns a Poisson-distributed unsigned int using Philox generator.
Definition rocrand_poisson.h:245
__forceinline__ __device__ __host__ uint4 rocrand_poisson4(rocrand_state_philox4x32_10 *state, double lambda)
Returns four Poisson-distributed unsigned int values using Philox generator.
Definition rocrand_poisson.h:264
__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:299
__forceinline__ __device__ __host__ double rocrand_normal_double(rocrand_state_philox4x32_10 *state)
Returns a normally distributed double value.
Definition rocrand_normal.h:442