/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-rocrand/checkouts/develop/projects/rocrand/library/include/rocrand/rocrand_poisson.h Source File

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-rocrand/checkouts/develop/projects/rocrand/library/include/rocrand/rocrand_poisson.h Source File#

API library: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-rocrand/checkouts/develop/projects/rocrand/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 
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 
51 namespace rocrand_device
52 {
53 namespace detail
54 {
55 
56 constexpr double lambda_threshold_small = 64.0;
57 constexpr double lambda_threshold_huge = 4000.0;
58 
59 template<class State, typename Result_Type = unsigned int>
60 __forceinline__ __device__ __host__
61 Result_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__
81 double 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 
112 template<class State, typename Result_Type = unsigned int>
113 __forceinline__ __device__ __host__
114 Result_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 
145 template<class State, typename Result_Type = unsigned int>
146 __forceinline__ __device__ __host__
147 Result_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 
155 template<class State, typename Result_Type = unsigned int>
156 __forceinline__ __device__ __host__
157 Result_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 
173 template<class State, typename Result_Type = unsigned int>
174 __forceinline__ __device__ __host__
175 Result_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 
215 template<class State, typename Result_Type = unsigned int>
216 __forceinline__ __device__ __host__
217 Result_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__
245 unsigned 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__
264 uint4 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__
295 unsigned 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__
316 unsigned 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__
337 unsigned 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__
357 unsigned 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__
376 unsigned 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__
395 unsigned 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__
413 unsigned 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__
432 unsigned 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__
450 unsigned 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__
469 unsigned 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__
486 unsigned 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__
503 unsigned 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__
520 unsigned 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