/build/reproducible-path/rocrand-6.4.1/library/include/rocrand/rocrand_normal.h Source File

/build/reproducible-path/rocrand-6.4.1/library/include/rocrand/rocrand_normal.h Source File#

API library: /build/reproducible-path/rocrand-6.4.1/library/include/rocrand/rocrand_normal.h Source File
rocrand_normal.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_NORMAL_H_
22#define ROCRAND_NORMAL_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_uniform.h"
47
48namespace rocrand_device {
49namespace detail {
50
51__forceinline__ __device__ __host__ float2 box_muller(unsigned int x, unsigned int y)
52{
53 float2 result;
54 float u = ROCRAND_2POW32_INV + (x * ROCRAND_2POW32_INV);
55 float v = ROCRAND_2POW32_INV_2PI + (y * ROCRAND_2POW32_INV_2PI);
56 float s = sqrtf(-2.0f * logf(u));
57 #ifdef __HIP_DEVICE_COMPILE__
58 __sincosf(v, &result.x, &result.y);
59 result.x *= s;
60 result.y *= s;
61 #else
62 result.x = sinf(v) * s;
63 result.y = cosf(v) * s;
64 #endif
65 return result;
66}
67
68__forceinline__ __device__ __host__ float2 box_muller(unsigned long long v)
69{
70 unsigned int x = static_cast<unsigned int>(v);
71 unsigned int y = static_cast<unsigned int>(v >> 32);
72
73 return box_muller(x, y);
74}
75
76__forceinline__ __device__ __host__ double2 box_muller_double(uint4 v)
77{
78 double2 result;
79 unsigned long long int v1 = (unsigned long long int)v.x ^
80 ((unsigned long long int)v.y << (53 - 32));
81 double u = ROCRAND_2POW53_INV_DOUBLE + (v1 * ROCRAND_2POW53_INV_DOUBLE);
82 unsigned long long int v2 = (unsigned long long int)v.z ^
83 ((unsigned long long int)v.w << (53 - 32));
84 double w = (ROCRAND_2POW53_INV_DOUBLE * 2.0) +
85 (v2 * (ROCRAND_2POW53_INV_DOUBLE * 2.0));
86 double s = sqrt(-2.0 * log(u));
87 #ifdef __HIP_DEVICE_COMPILE__
88 sincospi(w, &result.x, &result.y);
89 result.x *= s;
90 result.y *= s;
91 #else
92 result.x = sin(w * ROCRAND_PI_DOUBLE) * s;
93 result.y = cos(w * ROCRAND_PI_DOUBLE) * s;
94 #endif
95 return result;
96}
97
98__forceinline__ __device__ __host__ double2 box_muller_double(ulonglong2 v)
99{
100 unsigned int x = static_cast<unsigned int>(v.x);
101 unsigned int y = static_cast<unsigned int>(v.x >> 32);
102 unsigned int z = static_cast<unsigned int>(v.y);
103 unsigned int w = static_cast<unsigned int>(v.y >> 32);
104
105 return box_muller_double(make_uint4(x, y, z, w));
106}
107
108__forceinline__ __device__ __host__ __half2 box_muller_half(unsigned short x, unsigned short y)
109{
110 #if defined(ROCRAND_HALF_MATH_SUPPORTED)
111 __half u = __float2half(ROCRAND_2POW16_INV + (x * ROCRAND_2POW16_INV));
112 __half v = __float2half(ROCRAND_2POW16_INV_2PI + (y * ROCRAND_2POW16_INV_2PI));
113 __half s = hsqrt(__hmul(__float2half(-2.0f), hlog(u)));
114 return __half2 {
115 __hmul(hsin(v), s),
116 __hmul(hcos(v), s)
117 };
118 #else
119 float2 r;
120 float u = ROCRAND_2POW16_INV + (x * ROCRAND_2POW16_INV);
121 float v = ROCRAND_2POW16_INV_2PI + (y * ROCRAND_2POW16_INV_2PI);
122 float s = sqrtf(-2.0f * logf(u));
123 #ifdef __HIP_DEVICE_COMPILE__
124 __sincosf(v, &r.x, &r.y);
125 r.x *= s;
126 r.y *= s;
127 #else
128 r.x = sinf(v) * s;
129 r.y = cosf(v) * s;
130 #endif
131 return __half2 {
132 __float2half(r.x),
133 __float2half(r.y)
134 };
135 #endif
136}
137
138template<typename state_type>
139__forceinline__ __device__ __host__ float2 mrg_box_muller(unsigned int x, unsigned int y)
140{
141 float2 result;
142 float u = rocrand_device::detail::mrg_uniform_distribution<state_type>(x);
143 float v = rocrand_device::detail::mrg_uniform_distribution<state_type>(y) * ROCRAND_2PI;
144 float s = sqrtf(-2.0f * logf(u));
145 #ifdef __HIP_DEVICE_COMPILE__
146 __sincosf(v, &result.x, &result.y);
147 result.x *= s;
148 result.y *= s;
149 #else
150 result.x = sinf(v) * s;
151 result.y = cosf(v) * s;
152 #endif
153 return result;
154}
155
156template<typename state_type>
157__forceinline__ __device__ __host__ double2 mrg_box_muller_double(unsigned int x, unsigned int y)
158{
159 double2 result;
160 double u = rocrand_device::detail::mrg_uniform_distribution<state_type>(x);
161 double v = rocrand_device::detail::mrg_uniform_distribution<state_type>(y) * 2.0;
162 double s = sqrt(-2.0 * log(u));
163 #ifdef __HIP_DEVICE_COMPILE__
164 sincospi(v, &result.x, &result.y);
165 result.x *= s;
166 result.y *= s;
167 #else
168 result.x = sin(v * ROCRAND_PI_DOUBLE) * s;
169 result.y = cos(v * ROCRAND_PI_DOUBLE) * s;
170 #endif
171 return result;
172}
173
174__forceinline__ __device__ __host__ float roc_f_erfinv(float x)
175{
176 float tt1, tt2, lnx, sgn;
177 sgn = (x < 0.0f) ? -1.0f : 1.0f;
178
179 x = (1.0f - x) * (1.0f + x);
180 lnx = logf(x);
181
182 #ifdef __HIP_DEVICE_COMPILE__
183 if (isnan(lnx))
184 #else
185 if (std::isnan(lnx))
186 #endif
187 return 1.0f;
188 #ifdef __HIP_DEVICE_COMPILE__
189 else if (isinf(lnx))
190 #else
191 else if (std::isinf(lnx))
192 #endif
193 return 0.0f;
194
195 tt1 = 2.0f / (ROCRAND_PI * 0.147f) + 0.5f * lnx;
196 tt2 = 1.0f / (0.147f) * lnx;
197
198 return(sgn * sqrtf(-tt1 + sqrtf(tt1 * tt1 - tt2)));
199}
200
201__forceinline__ __device__ __host__ double roc_d_erfinv(double x)
202{
203 double tt1, tt2, lnx, sgn;
204 sgn = (x < 0.0) ? -1.0 : 1.0;
205
206 x = (1.0 - x) * (1.0 + x);
207 lnx = log(x);
208
209 #ifdef __HIP_DEVICE_COMPILE__
210 if (isnan(lnx))
211 #else
212 if (std::isnan(lnx))
213 #endif
214 return 1.0;
215 #ifdef __HIP_DEVICE_COMPILE__
216 else if (isinf(lnx))
217 #else
218 else if (std::isinf(lnx))
219 #endif
220 return 0.0;
221
222 tt1 = 2.0 / (ROCRAND_PI_DOUBLE * 0.147) + 0.5 * lnx;
223 tt2 = 1.0 / (0.147) * lnx;
224
225 return(sgn * sqrt(-tt1 + sqrt(tt1 * tt1 - tt2)));
226}
227
228__forceinline__ __device__ __host__ float normal_distribution(unsigned int x)
229{
230 float p = ::rocrand_device::detail::uniform_distribution(x);
231 float v = ROCRAND_SQRT2 * ::rocrand_device::detail::roc_f_erfinv(2.0f * p - 1.0f);
232 return v;
233}
234
235__forceinline__ __device__ __host__ float normal_distribution(unsigned long long int x)
236{
237 float p = ::rocrand_device::detail::uniform_distribution(x);
238 float v = ROCRAND_SQRT2 * ::rocrand_device::detail::roc_f_erfinv(2.0f * p - 1.0f);
239 return v;
240}
241
242__forceinline__ __device__ __host__ float2 normal_distribution2(unsigned int v1, unsigned int v2)
243{
244 return ::rocrand_device::detail::box_muller(v1, v2);
245}
246
247__forceinline__ __device__ __host__ float2 normal_distribution2(uint2 v)
248{
249 return ::rocrand_device::detail::box_muller(v.x, v.y);
250}
251
252__forceinline__ __device__ __host__ float2 normal_distribution2(unsigned long long v)
253{
254 return ::rocrand_device::detail::box_muller(v);
255}
256
257__forceinline__ __device__ __host__ float4 normal_distribution4(uint4 v)
258{
259 float2 r1 = ::rocrand_device::detail::box_muller(v.x, v.y);
260 float2 r2 = ::rocrand_device::detail::box_muller(v.z, v.w);
261 return float4{
262 r1.x,
263 r1.y,
264 r2.x,
265 r2.y
266 };
267}
268
269__forceinline__ __device__ __host__ float4 normal_distribution4(longlong2 v)
270{
271 float2 r1 = ::rocrand_device::detail::box_muller(v.x);
272 float2 r2 = ::rocrand_device::detail::box_muller(v.y);
273 return float4{r1.x, r1.y, r2.x, r2.y};
274}
275
276__forceinline__ __device__ __host__ float4 normal_distribution4(unsigned long long v1,
277 unsigned long long v2)
278{
279 float2 r1 = ::rocrand_device::detail::box_muller(v1);
280 float2 r2 = ::rocrand_device::detail::box_muller(v2);
281 return float4{r1.x, r1.y, r2.x, r2.y};
282}
283
284__forceinline__ __device__ __host__ double normal_distribution_double(unsigned int x)
285{
286 double p = ::rocrand_device::detail::uniform_distribution_double(x);
287 double v = ROCRAND_SQRT2 * ::rocrand_device::detail::roc_d_erfinv(2.0 * p - 1.0);
288 return v;
289}
290
291__forceinline__ __device__ __host__ double normal_distribution_double(unsigned long long int x)
292{
293 double p = ::rocrand_device::detail::uniform_distribution_double(x);
294 double v = ROCRAND_SQRT2 * ::rocrand_device::detail::roc_d_erfinv(2.0 * p - 1.0);
295 return v;
296}
297
298__forceinline__ __device__ __host__ double2 normal_distribution_double2(uint4 v)
299{
300 return ::rocrand_device::detail::box_muller_double(v);
301}
302
303__forceinline__ __device__ __host__ double2 normal_distribution_double2(ulonglong2 v)
304{
305 return ::rocrand_device::detail::box_muller_double(v);
306}
307
308__forceinline__ __device__ __host__ __half2 normal_distribution_half2(unsigned int v)
309{
310 return ::rocrand_device::detail::box_muller_half(
311 static_cast<unsigned short>(v),
312 static_cast<unsigned short>(v >> 16)
313 );
314}
315
316__forceinline__ __device__ __host__ __half2 normal_distribution_half2(unsigned long long v)
317{
318 return ::rocrand_device::detail::box_muller_half(static_cast<unsigned short>(v),
319 static_cast<unsigned short>(v >> 32));
320}
321
322template<typename state_type>
323__forceinline__ __device__ __host__ float2 mrg_normal_distribution2(unsigned int v1,
324 unsigned int v2)
325{
326 return ::rocrand_device::detail::mrg_box_muller<state_type>(v1, v2);
327}
328
329template<typename state_type>
330__forceinline__ __device__ __host__ double2 mrg_normal_distribution_double2(unsigned int v1,
331 unsigned int v2)
332{
333 return ::rocrand_device::detail::mrg_box_muller_double<state_type>(v1, v2);
334}
335
336template<typename state_type>
337__forceinline__ __device__ __host__ __half2 mrg_normal_distribution_half2(unsigned int v)
338{
339 v = rocrand_device::detail::mrg_uniform_distribution_uint<state_type>(v);
340 return ::rocrand_device::detail::box_muller_half(
341 static_cast<unsigned short>(v),
342 static_cast<unsigned short>(v >> 16)
343 );
344}
345
346} // end namespace detail
347} // end namespace rocrand_device
348
363#ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE
364__forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_philox4x32_10* state)
365{
366 typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_philox4x32_10> bm_helper;
367
368 if(bm_helper::has_float(state))
369 {
370 return bm_helper::get_float(state);
371 }
372
373 auto state1 = rocrand(state);
374 auto state2 = rocrand(state);
375
376 float2 r = rocrand_device::detail::normal_distribution2(state1, state2);
377 bm_helper::save_float(state, r.y);
378 return r.x;
379}
380#endif // ROCRAND_DETAIL_BM_NOT_IN_STATE
381
396__forceinline__ __device__ __host__ float2 rocrand_normal2(rocrand_state_philox4x32_10* state)
397{
398 auto state1 = rocrand(state);
399 auto state2 = rocrand(state);
400
401 return rocrand_device::detail::normal_distribution2(state1, state2);
402}
403
418__forceinline__ __device__ __host__ float4 rocrand_normal4(rocrand_state_philox4x32_10* state)
419{
420 return rocrand_device::detail::normal_distribution4(rocrand4(state));
421}
422
437#ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE
438__forceinline__ __device__ __host__ double rocrand_normal_double(rocrand_state_philox4x32_10* state)
439{
440 typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_philox4x32_10> bm_helper;
441
442 if(bm_helper::has_double(state))
443 {
444 return bm_helper::get_double(state);
445 }
446 double2 r = rocrand_device::detail::normal_distribution_double2(rocrand4(state));
447 bm_helper::save_double(state, r.y);
448 return r.x;
449}
450#endif // ROCRAND_DETAIL_BM_NOT_IN_STATE
451
466__forceinline__ __device__ __host__ double2
467 rocrand_normal_double2(rocrand_state_philox4x32_10* state)
468{
469 return rocrand_device::detail::normal_distribution_double2(rocrand4(state));
470}
471
486__forceinline__ __device__ __host__ double4
487 rocrand_normal_double4(rocrand_state_philox4x32_10* state)
488{
489 double2 r1, r2;
490 r1 = rocrand_device::detail::normal_distribution_double2(rocrand4(state));
491 r2 = rocrand_device::detail::normal_distribution_double2(rocrand4(state));
492 return double4 {
493 r1.x, r1.y, r2.x, r2.y
494 };
495}
496
511#ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE
512__forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_mrg31k3p* state)
513{
514 typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_mrg31k3p> bm_helper;
515
516 if(bm_helper::has_float(state))
517 {
518 return bm_helper::get_float(state);
519 }
520
521 auto state1 = state->next();
522 auto state2 = state->next();
523
524 float2 r
525 = rocrand_device::detail::mrg_normal_distribution2<rocrand_state_mrg31k3p>(state1, state2);
526 bm_helper::save_float(state, r.y);
527 return r.x;
528}
529#endif // ROCRAND_DETAIL_BM_NOT_IN_STATE
530
545__forceinline__ __device__ __host__ float2 rocrand_normal2(rocrand_state_mrg31k3p* state)
546{
547 auto state1 = state->next();
548 auto state2 = state->next();
549
550 return rocrand_device::detail::mrg_normal_distribution2<rocrand_state_mrg31k3p>(state1, state2);
551}
552
567#ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE
568__forceinline__ __device__ __host__ double rocrand_normal_double(rocrand_state_mrg31k3p* state)
569{
570 typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_mrg31k3p> bm_helper;
571
572 if(bm_helper::has_double(state))
573 {
574 return bm_helper::get_double(state);
575 }
576
577 auto state1 = state->next();
578 auto state2 = state->next();
579
580 double2 r
581 = rocrand_device::detail::mrg_normal_distribution_double2<rocrand_state_mrg31k3p>(state1,
582 state2);
583 bm_helper::save_double(state, r.y);
584 return r.x;
585}
586#endif // ROCRAND_DETAIL_BM_NOT_IN_STATE
587
602__forceinline__ __device__ __host__ double2 rocrand_normal_double2(rocrand_state_mrg31k3p* state)
603{
604 auto state1 = state->next();
605 auto state2 = state->next();
606
607 return rocrand_device::detail::mrg_normal_distribution_double2<rocrand_state_mrg31k3p>(state1,
608 state2);
609}
610
625#ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE
626__forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_mrg32k3a* state)
627{
628 typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_mrg32k3a> bm_helper;
629
630 if(bm_helper::has_float(state))
631 {
632 return bm_helper::get_float(state);
633 }
634
635 auto state1 = state->next();
636 auto state2 = state->next();
637
638 float2 r
639 = rocrand_device::detail::mrg_normal_distribution2<rocrand_state_mrg32k3a>(state1, state2);
640 bm_helper::save_float(state, r.y);
641 return r.x;
642}
643#endif // ROCRAND_DETAIL_BM_NOT_IN_STATE
644
659__forceinline__ __device__ __host__ float2 rocrand_normal2(rocrand_state_mrg32k3a* state)
660{
661 auto state1 = state->next();
662 auto state2 = state->next();
663
664 return rocrand_device::detail::mrg_normal_distribution2<rocrand_state_mrg32k3a>(state1, state2);
665}
666
681#ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE
682__forceinline__ __device__ __host__ double rocrand_normal_double(rocrand_state_mrg32k3a* state)
683{
684 typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_mrg32k3a> bm_helper;
685
686 if(bm_helper::has_double(state))
687 {
688 return bm_helper::get_double(state);
689 }
690
691 auto state1 = state->next();
692 auto state2 = state->next();
693
694 double2 r
695 = rocrand_device::detail::mrg_normal_distribution_double2<rocrand_state_mrg32k3a>(state1,
696 state2);
697 bm_helper::save_double(state, r.y);
698 return r.x;
699}
700#endif // ROCRAND_DETAIL_BM_NOT_IN_STATE
701
716__forceinline__ __device__ __host__ double2 rocrand_normal_double2(rocrand_state_mrg32k3a* state)
717{
718 auto state1 = state->next();
719 auto state2 = state->next();
720
721 return rocrand_device::detail::mrg_normal_distribution_double2<rocrand_state_mrg32k3a>(state1,
722 state2);
723}
724
739#ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE
740__forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_xorwow* state)
741{
742 typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_xorwow> bm_helper;
743
744 if(bm_helper::has_float(state))
745 {
746 return bm_helper::get_float(state);
747 }
748 auto state1 = rocrand(state);
749 auto state2 = rocrand(state);
750 float2 r = rocrand_device::detail::normal_distribution2(state1, state2);
751 bm_helper::save_float(state, r.y);
752 return r.x;
753}
754#endif // ROCRAND_DETAIL_BM_NOT_IN_STATE
755
770__forceinline__ __device__ __host__ float2 rocrand_normal2(rocrand_state_xorwow* state)
771{
772 auto state1 = rocrand(state);
773 auto state2 = rocrand(state);
774 return rocrand_device::detail::normal_distribution2(state1, state2);
775}
776
791#ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE
792__forceinline__ __device__ __host__ double rocrand_normal_double(rocrand_state_xorwow* state)
793{
794 typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_xorwow> bm_helper;
795
796 if(bm_helper::has_double(state))
797 {
798 return bm_helper::get_double(state);
799 }
800
801 auto state1 = rocrand(state);
802 auto state2 = rocrand(state);
803 auto state3 = rocrand(state);
804 auto state4 = rocrand(state);
805
806 double2 r = rocrand_device::detail::normal_distribution_double2(
807 uint4 { state1, state2, state3, state4 }
808 );
809 bm_helper::save_double(state, r.y);
810 return r.x;
811}
812#endif // ROCRAND_DETAIL_BM_NOT_IN_STATE
813
828__forceinline__ __device__ __host__ double2 rocrand_normal_double2(rocrand_state_xorwow* state)
829{
830 auto state1 = rocrand(state);
831 auto state2 = rocrand(state);
832 auto state3 = rocrand(state);
833 auto state4 = rocrand(state);
834
835 return rocrand_device::detail::normal_distribution_double2(
836 uint4 { state1, state2, state3, state4 }
837 );
838}
839
852__forceinline__ __device__ float rocrand_normal(rocrand_state_mtgp32* state)
853{
854 return rocrand_device::detail::normal_distribution(rocrand(state));
855}
856
871__forceinline__ __device__ float2 rocrand_normal2(rocrand_state_mtgp32* state)
872{
873 auto state1 = rocrand(state);
874 auto state2 = rocrand(state);
875 return rocrand_device::detail::normal_distribution2(state1, state2);
876}
877
890__forceinline__ __device__ double rocrand_normal_double(rocrand_state_mtgp32* state)
891{
892 return rocrand_device::detail::normal_distribution_double(rocrand(state));
893}
894
909__forceinline__ __device__ double2 rocrand_normal_double2(rocrand_state_mtgp32* state)
910{
911 auto state1 = rocrand(state);
912 auto state2 = rocrand(state);
913 auto state3 = rocrand(state);
914 auto state4 = rocrand(state);
915
916 return rocrand_device::detail::normal_distribution_double2(
917 uint4{state1, state2, state3, state4});
918}
919
932__forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_sobol32* state)
933{
934 return rocrand_device::detail::normal_distribution(rocrand(state));
935}
936
949__forceinline__ __device__ __host__ double rocrand_normal_double(rocrand_state_sobol32* state)
950{
951 return rocrand_device::detail::normal_distribution_double(rocrand(state));
952}
953
966__forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_scrambled_sobol32* state)
967{
968 return rocrand_device::detail::normal_distribution(rocrand(state));
969}
970
983__forceinline__ __device__ __host__ double
984 rocrand_normal_double(rocrand_state_scrambled_sobol32* state)
985{
986 return rocrand_device::detail::normal_distribution_double(rocrand(state));
987}
988
1001__forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_sobol64* state)
1002{
1003 return rocrand_device::detail::normal_distribution(rocrand(state));
1004}
1005
1018__forceinline__ __device__ __host__ double rocrand_normal_double(rocrand_state_sobol64* state)
1019{
1020 return rocrand_device::detail::normal_distribution_double(rocrand(state));
1021}
1022
1035__forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_scrambled_sobol64* state)
1036{
1037 return rocrand_device::detail::normal_distribution(rocrand(state));
1038}
1039
1052__forceinline__ __device__ __host__ double
1053 rocrand_normal_double(rocrand_state_scrambled_sobol64* state)
1054{
1055 return rocrand_device::detail::normal_distribution_double(rocrand(state));
1056}
1057
1070__forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_lfsr113* state)
1071{
1072 return rocrand_device::detail::normal_distribution(rocrand(state));
1073}
1074
1089__forceinline__ __device__ __host__ float2 rocrand_normal2(rocrand_state_lfsr113* state)
1090{
1091 auto state1 = rocrand(state);
1092 auto state2 = rocrand(state);
1093
1094 return rocrand_device::detail::normal_distribution2(state1, state2);
1095}
1096
1109__forceinline__ __device__ __host__ double rocrand_normal_double(rocrand_state_lfsr113* state)
1110{
1111 return rocrand_device::detail::normal_distribution_double(rocrand(state));
1112}
1113
1128__forceinline__ __device__ __host__ double2 rocrand_normal_double2(rocrand_state_lfsr113* state)
1129{
1130 auto state1 = rocrand(state);
1131 auto state2 = rocrand(state);
1132 auto state3 = rocrand(state);
1133 auto state4 = rocrand(state);
1134
1135 return rocrand_device::detail::normal_distribution_double2(
1136 uint4{state1, state2, state3, state4});
1137}
1138
1151__forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_threefry2x32_20* state)
1152{
1153 return rocrand_device::detail::normal_distribution(rocrand(state));
1154}
1155
1170__forceinline__ __device__ __host__ float2 rocrand_normal2(rocrand_state_threefry2x32_20* state)
1171{
1172 return rocrand_device::detail::normal_distribution2(rocrand2(state));
1173}
1174
1187__forceinline__ __device__ __host__ double
1188 rocrand_normal_double(rocrand_state_threefry2x32_20* state)
1189{
1190 return rocrand_device::detail::normal_distribution_double(rocrand(state));
1191}
1192
1207__forceinline__ __device__ __host__ double2
1208 rocrand_normal_double2(rocrand_state_threefry2x32_20* state)
1209{
1210 auto state1 = rocrand2(state);
1211 auto state2 = rocrand2(state);
1212
1213 return rocrand_device::detail::normal_distribution_double2(
1214 uint4{state1.x, state1.y, state2.x, state2.y});
1215}
1216
1229__forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_threefry2x64_20* state)
1230{
1231 return rocrand_device::detail::normal_distribution(rocrand(state));
1232}
1233
1248__forceinline__ __device__ __host__ float2 rocrand_normal2(rocrand_state_threefry2x64_20* state)
1249{
1250 return rocrand_device::detail::normal_distribution2(rocrand(state));
1251}
1252
1265__forceinline__ __device__ __host__ double
1266 rocrand_normal_double(rocrand_state_threefry2x64_20* state)
1267{
1268 return rocrand_device::detail::normal_distribution_double(rocrand(state));
1269}
1270
1285__forceinline__ __device__ __host__ double2
1286 rocrand_normal_double2(rocrand_state_threefry2x64_20* state)
1287{
1288 return rocrand_device::detail::normal_distribution_double2(rocrand2(state));
1289}
1290
1303__forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_threefry4x32_20* state)
1304{
1305 return rocrand_device::detail::normal_distribution(rocrand(state));
1306}
1307
1322__forceinline__ __device__ __host__ float2 rocrand_normal2(rocrand_state_threefry4x32_20* state)
1323{
1324 auto state1 = rocrand(state);
1325 auto state2 = rocrand(state);
1326
1327 return rocrand_device::detail::normal_distribution2(state1, state2);
1328}
1329
1342__forceinline__ __device__ __host__ double
1343 rocrand_normal_double(rocrand_state_threefry4x32_20* state)
1344{
1345 return rocrand_device::detail::normal_distribution_double(rocrand(state));
1346}
1347
1362__forceinline__ __device__ __host__ double2
1363 rocrand_normal_double2(rocrand_state_threefry4x32_20* state)
1364{
1365 return rocrand_device::detail::normal_distribution_double2(rocrand4(state));
1366}
1367
1380__forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_threefry4x64_20* state)
1381{
1382 return rocrand_device::detail::normal_distribution(rocrand(state));
1383}
1384
1399__forceinline__ __device__ __host__ float2 rocrand_normal2(rocrand_state_threefry4x64_20* state)
1400{
1401 auto state1 = rocrand(state);
1402 auto state2 = rocrand(state);
1403
1404 return rocrand_device::detail::normal_distribution2(state1, state2);
1405}
1406
1419__forceinline__ __device__ __host__ double
1420 rocrand_normal_double(rocrand_state_threefry4x64_20* state)
1421{
1422 return rocrand_device::detail::normal_distribution_double(rocrand(state));
1423}
1424
1439__forceinline__ __device__ __host__ double2
1440 rocrand_normal_double2(rocrand_state_threefry4x64_20* state)
1441{
1442 auto state1 = rocrand(state);
1443 auto state2 = rocrand(state);
1444
1445 return rocrand_device::detail::normal_distribution_double2(ulonglong2{state1, state2});
1446}
1447
// end of group rocranddevice
1449
1450#endif // ROCRAND_NORMAL_H_
__forceinline__ __device__ __host__ double4 rocrand_normal_double4(rocrand_state_philox4x32_10 *state)
Returns four normally distributed double values.
Definition rocrand_normal.h:487
__forceinline__ __device__ __host__ double2 rocrand_normal_double2(rocrand_state_philox4x32_10 *state)
Returns two normally distributed double values.
Definition rocrand_normal.h:467
__forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_philox4x32_10 *state)
Returns a normally distributed float value.
Definition rocrand_normal.h:364
__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__ double rocrand_normal_double(rocrand_state_philox4x32_10 *state)
Returns a normally distributed double value.
Definition rocrand_normal.h:438
__forceinline__ __device__ __host__ float4 rocrand_normal4(rocrand_state_philox4x32_10 *state)
Returns four normally distributed float values.
Definition rocrand_normal.h:418
__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
__forceinline__ __device__ __host__ float2 rocrand_normal2(rocrand_state_philox4x32_10 *state)
Returns two normally distributed float values.
Definition rocrand_normal.h:396