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

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

API library: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-rocrand/checkouts/latest/library/include/rocrand/rocrand_normal.h Source File
API library
rocrand_normal.h
1 // Copyright (c) 2017-2022 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_NORMAL_H_
22 #define ROCRAND_NORMAL_H_
23 
24 #ifndef FQUALIFIERS
25 #define FQUALIFIERS __forceinline__ __device__
26 #endif // FQUALIFIERS
27 
33 #include <math.h>
34 
35 #include "rocrand/rocrand_lfsr113.h"
36 #include "rocrand/rocrand_mrg31k3p.h"
37 #include "rocrand/rocrand_mrg32k3a.h"
38 #include "rocrand/rocrand_mtgp32.h"
39 #include "rocrand/rocrand_philox4x32_10.h"
40 #include "rocrand/rocrand_scrambled_sobol32.h"
41 #include "rocrand/rocrand_scrambled_sobol64.h"
42 #include "rocrand/rocrand_sobol32.h"
43 #include "rocrand/rocrand_sobol64.h"
44 #include "rocrand/rocrand_threefry2x32_20.h"
45 #include "rocrand/rocrand_threefry2x64_20.h"
46 #include "rocrand/rocrand_threefry4x32_20.h"
47 #include "rocrand/rocrand_threefry4x64_20.h"
48 #include "rocrand/rocrand_xorwow.h"
49 
50 #include "rocrand/rocrand_uniform.h"
51 
52 namespace rocrand_device {
53 namespace detail {
54 
56 float2 box_muller(unsigned int x, unsigned int y)
57 {
58  float2 result;
59  float u = ROCRAND_2POW32_INV + (x * ROCRAND_2POW32_INV);
60  float v = ROCRAND_2POW32_INV_2PI + (y * ROCRAND_2POW32_INV_2PI);
61  float s = sqrtf(-2.0f * logf(u));
62  #ifdef __HIP_DEVICE_COMPILE__
63  __sincosf(v, &result.x, &result.y);
64  result.x *= s;
65  result.y *= s;
66  #else
67  result.x = sinf(v) * s;
68  result.y = cosf(v) * s;
69  #endif
70  return result;
71 }
72 
73 FQUALIFIERS float2 box_muller(unsigned long long v)
74 {
75  unsigned int x = static_cast<unsigned int>(v);
76  unsigned int y = static_cast<unsigned int>(v >> 32);
77 
78  return box_muller(x, y);
79 }
80 
82 double2 box_muller_double(uint4 v)
83 {
84  double2 result;
85  unsigned long long int v1 = (unsigned long long int)v.x ^
86  ((unsigned long long int)v.y << (53 - 32));
87  double u = ROCRAND_2POW53_INV_DOUBLE + (v1 * ROCRAND_2POW53_INV_DOUBLE);
88  unsigned long long int v2 = (unsigned long long int)v.z ^
89  ((unsigned long long int)v.w << (53 - 32));
90  double w = (ROCRAND_2POW53_INV_DOUBLE * 2.0) +
91  (v2 * (ROCRAND_2POW53_INV_DOUBLE * 2.0));
92  double s = sqrt(-2.0 * log(u));
93  #ifdef __HIP_DEVICE_COMPILE__
94  sincospi(w, &result.x, &result.y);
95  result.x *= s;
96  result.y *= s;
97  #else
98  result.x = sin(w * ROCRAND_PI_DOUBLE) * s;
99  result.y = cos(w * ROCRAND_PI_DOUBLE) * s;
100  #endif
101  return result;
102 }
103 
104 FQUALIFIERS double2 box_muller_double(ulonglong2 v)
105 {
106  unsigned int x = static_cast<unsigned int>(v.x);
107  unsigned int y = static_cast<unsigned int>(v.x >> 32);
108  unsigned int z = static_cast<unsigned int>(v.y);
109  unsigned int w = static_cast<unsigned int>(v.y >> 32);
110 
111  return box_muller_double(make_uint4(x, y, z, w));
112 }
113 
115 __half2 box_muller_half(unsigned short x, unsigned short y)
116 {
117  #if defined(ROCRAND_HALF_MATH_SUPPORTED)
118  __half u = __float2half(ROCRAND_2POW16_INV + (x * ROCRAND_2POW16_INV));
119  __half v = __float2half(ROCRAND_2POW16_INV_2PI + (y * ROCRAND_2POW16_INV_2PI));
120  __half s = hsqrt(__hmul(__float2half(-2.0f), hlog(u)));
121  return __half2 {
122  __hmul(hsin(v), s),
123  __hmul(hcos(v), s)
124  };
125  #else
126  float2 r;
127  float u = ROCRAND_2POW16_INV + (x * ROCRAND_2POW16_INV);
128  float v = ROCRAND_2POW16_INV_2PI + (y * ROCRAND_2POW16_INV_2PI);
129  float s = sqrtf(-2.0f * logf(u));
130  #ifdef __HIP_DEVICE_COMPILE__
131  __sincosf(v, &r.x, &r.y);
132  r.x *= s;
133  r.y *= s;
134  #else
135  r.x = sinf(v) * s;
136  r.y = cosf(v) * s;
137  #endif
138  return __half2 {
139  __float2half(r.x),
140  __float2half(r.y)
141  };
142  #endif
143 }
144 
145 template<typename state_type>
146 FQUALIFIERS float2 mrg_box_muller(unsigned int x, unsigned int y)
147 {
148  float2 result;
149  float u = rocrand_device::detail::mrg_uniform_distribution<state_type>(x);
150  float v = rocrand_device::detail::mrg_uniform_distribution<state_type>(y) * ROCRAND_2PI;
151  float s = sqrtf(-2.0f * logf(u));
152  #ifdef __HIP_DEVICE_COMPILE__
153  __sincosf(v, &result.x, &result.y);
154  result.x *= s;
155  result.y *= s;
156  #else
157  result.x = sinf(v) * s;
158  result.y = cosf(v) * s;
159  #endif
160  return result;
161 }
162 
163 template<typename state_type>
164 FQUALIFIERS double2 mrg_box_muller_double(unsigned int x, unsigned int y)
165 {
166  double2 result;
167  double u = rocrand_device::detail::mrg_uniform_distribution<state_type>(x);
168  double v = rocrand_device::detail::mrg_uniform_distribution<state_type>(y) * 2.0;
169  double s = sqrt(-2.0 * log(u));
170  #ifdef __HIP_DEVICE_COMPILE__
171  sincospi(v, &result.x, &result.y);
172  result.x *= s;
173  result.y *= s;
174  #else
175  result.x = sin(v * ROCRAND_PI_DOUBLE) * s;
176  result.y = cos(v * ROCRAND_PI_DOUBLE) * s;
177  #endif
178  return result;
179 }
180 
182 float roc_f_erfinv(float x)
183 {
184  float tt1, tt2, lnx, sgn;
185  sgn = (x < 0.0f) ? -1.0f : 1.0f;
186 
187  x = (1.0f - x) * (1.0f + x);
188  lnx = logf(x);
189 
190  #ifdef __HIP_DEVICE_COMPILE__
191  if (isnan(lnx))
192  #else
193  if (std::isnan(lnx))
194  #endif
195  return 1.0f;
196  #ifdef __HIP_DEVICE_COMPILE__
197  else if (isinf(lnx))
198  #else
199  else if (std::isinf(lnx))
200  #endif
201  return 0.0f;
202 
203  tt1 = 2.0f / (ROCRAND_PI * 0.147f) + 0.5f * lnx;
204  tt2 = 1.0f / (0.147f) * lnx;
205 
206  return(sgn * sqrtf(-tt1 + sqrtf(tt1 * tt1 - tt2)));
207 }
208 
210 double roc_d_erfinv(double x)
211 {
212  double tt1, tt2, lnx, sgn;
213  sgn = (x < 0.0) ? -1.0 : 1.0;
214 
215  x = (1.0 - x) * (1.0 + x);
216  lnx = log(x);
217 
218  #ifdef __HIP_DEVICE_COMPILE__
219  if (isnan(lnx))
220  #else
221  if (std::isnan(lnx))
222  #endif
223  return 1.0;
224  #ifdef __HIP_DEVICE_COMPILE__
225  else if (isinf(lnx))
226  #else
227  else if (std::isinf(lnx))
228  #endif
229  return 0.0;
230 
231  tt1 = 2.0 / (ROCRAND_PI_DOUBLE * 0.147) + 0.5 * lnx;
232  tt2 = 1.0 / (0.147) * lnx;
233 
234  return(sgn * sqrt(-tt1 + sqrt(tt1 * tt1 - tt2)));
235 }
236 
238 float normal_distribution(unsigned int x)
239 {
240  float p = ::rocrand_device::detail::uniform_distribution(x);
241  float v = ROCRAND_SQRT2 * ::rocrand_device::detail::roc_f_erfinv(2.0f * p - 1.0f);
242  return v;
243 }
244 
246 float normal_distribution(unsigned long long int x)
247 {
248  float p = ::rocrand_device::detail::uniform_distribution(x);
249  float v = ROCRAND_SQRT2 * ::rocrand_device::detail::roc_f_erfinv(2.0f * p - 1.0f);
250  return v;
251 }
252 
254 float2 normal_distribution2(unsigned int v1, unsigned int v2)
255 {
256  return ::rocrand_device::detail::box_muller(v1, v2);
257 }
258 
259 FQUALIFIERS float2 normal_distribution2(uint2 v)
260 {
261  return ::rocrand_device::detail::box_muller(v.x, v.y);
262 }
263 
264 FQUALIFIERS float2 normal_distribution2(unsigned long long v)
265 {
266  return ::rocrand_device::detail::box_muller(v);
267 }
268 
270 float4 normal_distribution4(uint4 v)
271 {
272  float2 r1 = ::rocrand_device::detail::box_muller(v.x, v.y);
273  float2 r2 = ::rocrand_device::detail::box_muller(v.z, v.w);
274  return float4{
275  r1.x,
276  r1.y,
277  r2.x,
278  r2.y
279  };
280 }
281 
282 FQUALIFIERS float4 normal_distribution4(longlong2 v)
283 {
284  float2 r1 = ::rocrand_device::detail::box_muller(v.x);
285  float2 r2 = ::rocrand_device::detail::box_muller(v.y);
286  return float4{r1.x, r1.y, r2.x, r2.y};
287 }
288 
289 FQUALIFIERS float4 normal_distribution4(unsigned long long v1, unsigned long long v2)
290 {
291  float2 r1 = ::rocrand_device::detail::box_muller(v1);
292  float2 r2 = ::rocrand_device::detail::box_muller(v2);
293  return float4{r1.x, r1.y, r2.x, r2.y};
294 }
295 
297 double normal_distribution_double(unsigned int x)
298 {
299  double p = ::rocrand_device::detail::uniform_distribution_double(x);
300  double v = ROCRAND_SQRT2 * ::rocrand_device::detail::roc_d_erfinv(2.0 * p - 1.0);
301  return v;
302 }
303 
305 double normal_distribution_double(unsigned long long int x)
306 {
307  double p = ::rocrand_device::detail::uniform_distribution_double(x);
308  double v = ROCRAND_SQRT2 * ::rocrand_device::detail::roc_d_erfinv(2.0 * p - 1.0);
309  return v;
310 }
311 
313 double2 normal_distribution_double2(uint4 v)
314 {
315  return ::rocrand_device::detail::box_muller_double(v);
316 }
317 
318 FQUALIFIERS double2 normal_distribution_double2(ulonglong2 v)
319 {
320  return ::rocrand_device::detail::box_muller_double(v);
321 }
322 
324 __half2 normal_distribution_half2(unsigned int v)
325 {
326  return ::rocrand_device::detail::box_muller_half(
327  static_cast<unsigned short>(v),
328  static_cast<unsigned short>(v >> 16)
329  );
330 }
331 
332 FQUALIFIERS __half2 normal_distribution_half2(unsigned long long v)
333 {
334  return ::rocrand_device::detail::box_muller_half(static_cast<unsigned short>(v),
335  static_cast<unsigned short>(v >> 32));
336 }
337 
338 template<typename state_type>
339 FQUALIFIERS float2 mrg_normal_distribution2(unsigned int v1, unsigned int v2)
340 {
341  return ::rocrand_device::detail::mrg_box_muller<state_type>(v1, v2);
342 }
343 
344 template<typename state_type>
345 FQUALIFIERS double2 mrg_normal_distribution_double2(unsigned int v1, unsigned int v2)
346 {
347  return ::rocrand_device::detail::mrg_box_muller_double<state_type>(v1, v2);
348 }
349 
350 template<typename state_type>
351 FQUALIFIERS __half2 mrg_normal_distribution_half2(unsigned int v)
352 {
353  v = rocrand_device::detail::mrg_uniform_distribution_uint<state_type>(v);
354  return ::rocrand_device::detail::box_muller_half(
355  static_cast<unsigned short>(v),
356  static_cast<unsigned short>(v >> 16)
357  );
358 }
359 
360 } // end namespace detail
361 } // end namespace rocrand_device
362 
377 #ifndef ROCRAND_DETAIL_PHILOX_BM_NOT_IN_STATE
379 float rocrand_normal(rocrand_state_philox4x32_10 * state)
380 {
381  typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_philox4x32_10> bm_helper;
382 
383  if(bm_helper::has_float(state))
384  {
385  return bm_helper::get_float(state);
386  }
387 
388  auto state1 = rocrand(state);
389  auto state2 = rocrand(state);
390 
391  float2 r = rocrand_device::detail::normal_distribution2(state1, state2);
392  bm_helper::save_float(state, r.y);
393  return r.x;
394 }
395 #endif // ROCRAND_DETAIL_PHILOX_BM_NOT_IN_STATE
396 
412 float2 rocrand_normal2(rocrand_state_philox4x32_10 * state)
413 {
414  auto state1 = rocrand(state);
415  auto state2 = rocrand(state);
416 
417  return rocrand_device::detail::normal_distribution2(state1, state2);
418 }
419 
435 float4 rocrand_normal4(rocrand_state_philox4x32_10 * state)
436 {
437  return rocrand_device::detail::normal_distribution4(rocrand4(state));
438 }
439 
454 #ifndef ROCRAND_DETAIL_PHILOX_BM_NOT_IN_STATE
456 double rocrand_normal_double(rocrand_state_philox4x32_10 * state)
457 {
458  typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_philox4x32_10> bm_helper;
459 
460  if(bm_helper::has_double(state))
461  {
462  return bm_helper::get_double(state);
463  }
464  double2 r = rocrand_device::detail::normal_distribution_double2(rocrand4(state));
465  bm_helper::save_double(state, r.y);
466  return r.x;
467 }
468 #endif // ROCRAND_DETAIL_PHILOX_BM_NOT_IN_STATE
469 
485 double2 rocrand_normal_double2(rocrand_state_philox4x32_10 * state)
486 {
487  return rocrand_device::detail::normal_distribution_double2(rocrand4(state));
488 }
489 
505 double4 rocrand_normal_double4(rocrand_state_philox4x32_10 * state)
506 {
507  double2 r1, r2;
508  r1 = rocrand_device::detail::normal_distribution_double2(rocrand4(state));
509  r2 = rocrand_device::detail::normal_distribution_double2(rocrand4(state));
510  return double4 {
511  r1.x, r1.y, r2.x, r2.y
512  };
513 }
514 
529 #ifndef ROCRAND_DETAIL_MRG31K3P_BM_NOT_IN_STATE
530 FQUALIFIERS float rocrand_normal(rocrand_state_mrg31k3p* state)
531 {
532  typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_mrg31k3p> bm_helper;
533 
534  if(bm_helper::has_float(state))
535  {
536  return bm_helper::get_float(state);
537  }
538 
539  auto state1 = state->next();
540  auto state2 = state->next();
541 
542  float2 r
543  = rocrand_device::detail::mrg_normal_distribution2<rocrand_state_mrg31k3p>(state1, state2);
544  bm_helper::save_float(state, r.y);
545  return r.x;
546 }
547 #endif // ROCRAND_DETAIL_MRG31K3P_BM_NOT_IN_STATE
548 
563 FQUALIFIERS float2 rocrand_normal2(rocrand_state_mrg31k3p* state)
564 {
565  auto state1 = state->next();
566  auto state2 = state->next();
567 
568  return rocrand_device::detail::mrg_normal_distribution2<rocrand_state_mrg31k3p>(state1, state2);
569 }
570 
585 #ifndef ROCRAND_DETAIL_MRG31K3P_BM_NOT_IN_STATE
586 FQUALIFIERS double rocrand_normal_double(rocrand_state_mrg31k3p* state)
587 {
588  typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_mrg31k3p> bm_helper;
589 
590  if(bm_helper::has_double(state))
591  {
592  return bm_helper::get_double(state);
593  }
594 
595  auto state1 = state->next();
596  auto state2 = state->next();
597 
598  double2 r
599  = rocrand_device::detail::mrg_normal_distribution_double2<rocrand_state_mrg31k3p>(state1,
600  state2);
601  bm_helper::save_double(state, r.y);
602  return r.x;
603 }
604 #endif // ROCRAND_DETAIL_MRG31K3P_BM_NOT_IN_STATE
605 
620 FQUALIFIERS double2 rocrand_normal_double2(rocrand_state_mrg31k3p* state)
621 {
622  auto state1 = state->next();
623  auto state2 = state->next();
624 
625  return rocrand_device::detail::mrg_normal_distribution_double2<rocrand_state_mrg31k3p>(state1,
626  state2);
627 }
628 
643 #ifndef ROCRAND_DETAIL_MRG32K3A_BM_NOT_IN_STATE
645 float rocrand_normal(rocrand_state_mrg32k3a * state)
646 {
647  typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_mrg32k3a> bm_helper;
648 
649  if(bm_helper::has_float(state))
650  {
651  return bm_helper::get_float(state);
652  }
653 
654  auto state1 = state->next();
655  auto state2 = state->next();
656 
657  float2 r
658  = rocrand_device::detail::mrg_normal_distribution2<rocrand_state_mrg32k3a>(state1, state2);
659  bm_helper::save_float(state, r.y);
660  return r.x;
661 }
662 #endif // ROCRAND_DETAIL_MRG32K3A_BM_NOT_IN_STATE
663 
679 float2 rocrand_normal2(rocrand_state_mrg32k3a * state)
680 {
681  auto state1 = state->next();
682  auto state2 = state->next();
683 
684  return rocrand_device::detail::mrg_normal_distribution2<rocrand_state_mrg32k3a>(state1, state2);
685 }
686 
701 #ifndef ROCRAND_DETAIL_MRG32K3A_BM_NOT_IN_STATE
703 double rocrand_normal_double(rocrand_state_mrg32k3a * state)
704 {
705  typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_mrg32k3a> bm_helper;
706 
707  if(bm_helper::has_double(state))
708  {
709  return bm_helper::get_double(state);
710  }
711 
712  auto state1 = state->next();
713  auto state2 = state->next();
714 
715  double2 r
716  = rocrand_device::detail::mrg_normal_distribution_double2<rocrand_state_mrg32k3a>(state1,
717  state2);
718  bm_helper::save_double(state, r.y);
719  return r.x;
720 }
721 #endif // ROCRAND_DETAIL_MRG32K3A_BM_NOT_IN_STATE
722 
738 double2 rocrand_normal_double2(rocrand_state_mrg32k3a * state)
739 {
740  auto state1 = state->next();
741  auto state2 = state->next();
742 
743  return rocrand_device::detail::mrg_normal_distribution_double2<rocrand_state_mrg32k3a>(state1,
744  state2);
745 }
746 
761 #ifndef ROCRAND_DETAIL_XORWOW_BM_NOT_IN_STATE
763 float rocrand_normal(rocrand_state_xorwow * state)
764 {
765  typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_xorwow> bm_helper;
766 
767  if(bm_helper::has_float(state))
768  {
769  return bm_helper::get_float(state);
770  }
771  auto state1 = rocrand(state);
772  auto state2 = rocrand(state);
773  float2 r = rocrand_device::detail::normal_distribution2(state1, state2);
774  bm_helper::save_float(state, r.y);
775  return r.x;
776 }
777 #endif // ROCRAND_DETAIL_XORWOW_BM_NOT_IN_STATE
778 
794 float2 rocrand_normal2(rocrand_state_xorwow * state)
795 {
796  auto state1 = rocrand(state);
797  auto state2 = rocrand(state);
798  return rocrand_device::detail::normal_distribution2(state1, state2);
799 }
800 
815 #ifndef ROCRAND_DETAIL_XORWOW_BM_NOT_IN_STATE
817 double rocrand_normal_double(rocrand_state_xorwow * state)
818 {
819  typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_xorwow> bm_helper;
820 
821  if(bm_helper::has_double(state))
822  {
823  return bm_helper::get_double(state);
824  }
825 
826  auto state1 = rocrand(state);
827  auto state2 = rocrand(state);
828  auto state3 = rocrand(state);
829  auto state4 = rocrand(state);
830 
831  double2 r = rocrand_device::detail::normal_distribution_double2(
832  uint4 { state1, state2, state3, state4 }
833  );
834  bm_helper::save_double(state, r.y);
835  return r.x;
836 }
837 #endif // ROCRAND_DETAIL_XORWOW_BM_NOT_IN_STATE
838 
854 double2 rocrand_normal_double2(rocrand_state_xorwow * state)
855 {
856  auto state1 = rocrand(state);
857  auto state2 = rocrand(state);
858  auto state3 = rocrand(state);
859  auto state4 = rocrand(state);
860 
861  return rocrand_device::detail::normal_distribution_double2(
862  uint4 { state1, state2, state3, state4 }
863  );
864 }
865 
879 float rocrand_normal(rocrand_state_mtgp32 * state)
880 {
881  return rocrand_device::detail::normal_distribution(rocrand(state));
882 }
883 
897 double rocrand_normal_double(rocrand_state_mtgp32 * state)
898 {
899  return rocrand_device::detail::normal_distribution_double(rocrand(state));
900 }
901 
915 float rocrand_normal(rocrand_state_sobol32 * state)
916 {
917  return rocrand_device::detail::normal_distribution(rocrand(state));
918 }
919 
933 double rocrand_normal_double(rocrand_state_sobol32 * state)
934 {
935  return rocrand_device::detail::normal_distribution_double(rocrand(state));
936 }
937 
951 float rocrand_normal(rocrand_state_scrambled_sobol32* state)
952 {
953  return rocrand_device::detail::normal_distribution(rocrand(state));
954 }
955 
969 double rocrand_normal_double(rocrand_state_scrambled_sobol32* state)
970 {
971  return rocrand_device::detail::normal_distribution_double(rocrand(state));
972 }
973 
987 float rocrand_normal(rocrand_state_sobol64* state)
988 {
989  return rocrand_device::detail::normal_distribution(rocrand(state));
990 }
991 
1005 double rocrand_normal_double(rocrand_state_sobol64 * state)
1006 {
1007  return rocrand_device::detail::normal_distribution_double(rocrand(state));
1008 }
1009 
1023 float rocrand_normal(rocrand_state_scrambled_sobol64* state)
1024 {
1025  return rocrand_device::detail::normal_distribution(rocrand(state));
1026 }
1027 
1041 double rocrand_normal_double(rocrand_state_scrambled_sobol64* state)
1042 {
1043  return rocrand_device::detail::normal_distribution_double(rocrand(state));
1044 }
1045 
1059 float rocrand_normal(rocrand_state_lfsr113* state)
1060 {
1061  return rocrand_device::detail::normal_distribution(rocrand(state));
1062 }
1063 
1079 float2 rocrand_normal2(rocrand_state_lfsr113* state)
1080 {
1081  auto state1 = rocrand(state);
1082  auto state2 = rocrand(state);
1083 
1084  return rocrand_device::detail::normal_distribution2(state1, state2);
1085 }
1086 
1100 double rocrand_normal_double(rocrand_state_lfsr113* state)
1101 {
1102  return rocrand_device::detail::normal_distribution_double(rocrand(state));
1103 }
1104 
1120 double2 rocrand_normal_double2(rocrand_state_lfsr113* state)
1121 {
1122  auto state1 = rocrand(state);
1123  auto state2 = rocrand(state);
1124  auto state3 = rocrand(state);
1125  auto state4 = rocrand(state);
1126 
1127  return rocrand_device::detail::normal_distribution_double2(
1128  uint4{state1, state2, state3, state4});
1129 }
1130 
1143 FQUALIFIERS float rocrand_normal(rocrand_state_threefry2x32_20* state)
1144 {
1145  return rocrand_device::detail::normal_distribution(rocrand(state));
1146 }
1147 
1162 FQUALIFIERS float2 rocrand_normal2(rocrand_state_threefry2x32_20* state)
1163 {
1164  return rocrand_device::detail::normal_distribution2(rocrand2(state));
1165 }
1166 
1179 FQUALIFIERS double rocrand_normal_double(rocrand_state_threefry2x32_20* state)
1180 {
1181  return rocrand_device::detail::normal_distribution_double(rocrand(state));
1182 }
1183 
1198 FQUALIFIERS double2 rocrand_normal_double2(rocrand_state_threefry2x32_20* state)
1199 {
1200  auto state1 = rocrand2(state);
1201  auto state2 = rocrand2(state);
1202 
1203  return rocrand_device::detail::normal_distribution_double2(
1204  uint4{state1.x, state1.y, state2.x, state2.y});
1205 }
1206 
1219 FQUALIFIERS float rocrand_normal(rocrand_state_threefry2x64_20* state)
1220 {
1221  return rocrand_device::detail::normal_distribution(rocrand(state));
1222 }
1223 
1238 FQUALIFIERS float2 rocrand_normal2(rocrand_state_threefry2x64_20* state)
1239 {
1240  return rocrand_device::detail::normal_distribution2(rocrand(state));
1241 }
1242 
1255 FQUALIFIERS double rocrand_normal_double(rocrand_state_threefry2x64_20* state)
1256 {
1257  return rocrand_device::detail::normal_distribution_double(rocrand(state));
1258 }
1259 
1274 FQUALIFIERS double2 rocrand_normal_double2(rocrand_state_threefry2x64_20* state)
1275 {
1276  return rocrand_device::detail::normal_distribution_double2(rocrand2(state));
1277 }
1278 
1291 FQUALIFIERS float rocrand_normal(rocrand_state_threefry4x32_20* state)
1292 {
1293  return rocrand_device::detail::normal_distribution(rocrand(state));
1294 }
1295 
1310 FQUALIFIERS float2 rocrand_normal2(rocrand_state_threefry4x32_20* state)
1311 {
1312  auto state1 = rocrand(state);
1313  auto state2 = rocrand(state);
1314 
1315  return rocrand_device::detail::normal_distribution2(state1, state2);
1316 }
1317 
1330 FQUALIFIERS double rocrand_normal_double(rocrand_state_threefry4x32_20* state)
1331 {
1332  return rocrand_device::detail::normal_distribution_double(rocrand(state));
1333 }
1334 
1349 FQUALIFIERS double2 rocrand_normal_double2(rocrand_state_threefry4x32_20* state)
1350 {
1351  return rocrand_device::detail::normal_distribution_double2(rocrand4(state));
1352 }
1353 
1366 FQUALIFIERS float rocrand_normal(rocrand_state_threefry4x64_20* state)
1367 {
1368  return rocrand_device::detail::normal_distribution(rocrand(state));
1369 }
1370 
1385 FQUALIFIERS float2 rocrand_normal2(rocrand_state_threefry4x64_20* state)
1386 {
1387  auto state1 = rocrand(state);
1388  auto state2 = rocrand(state);
1389 
1390  return rocrand_device::detail::normal_distribution2(state1, state2);
1391 }
1392 
1405 FQUALIFIERS double rocrand_normal_double(rocrand_state_threefry4x64_20* state)
1406 {
1407  return rocrand_device::detail::normal_distribution_double(rocrand(state));
1408 }
1409 
1424 FQUALIFIERS double2 rocrand_normal_double2(rocrand_state_threefry4x64_20* state)
1425 {
1426  auto state1 = rocrand(state);
1427  auto state2 = rocrand(state);
1428 
1429  return rocrand_device::detail::normal_distribution_double2(ulonglong2{state1, state2});
1430 }
1431  // end of group rocranddevice
1433 
1434 #endif // ROCRAND_NORMAL_H_
FQUALIFIERS float2 rocrand_normal2(rocrand_state_philox4x32_10 *state)
Returns two normally distributed float values.
Definition: rocrand_normal.h:412
FQUALIFIERS double rocrand_normal_double(rocrand_state_philox4x32_10 *state)
Returns a normally distributed double value.
Definition: rocrand_normal.h:456
FQUALIFIERS double4 rocrand_normal_double4(rocrand_state_philox4x32_10 *state)
Returns four normally distributed double values.
Definition: rocrand_normal.h:505
FQUALIFIERS unsigned int rocrand(rocrand_state_lfsr113 *state)
Returns uniformly distributed random unsigned int value from [0; 2^32 - 1] range.
Definition: rocrand_lfsr113.h:253
FQUALIFIERS float4 rocrand_normal4(rocrand_state_philox4x32_10 *state)
Returns four normally distributed float values.
Definition: rocrand_normal.h:435
FQUALIFIERS float rocrand_normal(rocrand_state_philox4x32_10 *state)
Returns a normally distributed float value.
Definition: rocrand_normal.h:379
FQUALIFIERS 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:410
FQUALIFIERS double2 rocrand_normal_double2(rocrand_state_philox4x32_10 *state)
Returns two normally distributed double values.
Definition: rocrand_normal.h:485
#define FQUALIFIERS
Shorthand for commonly used function qualifiers.
Definition: rocrand_uniform.h:31