Skip to main content
`Ctrl`+`K`

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

API library: /build/rocrand-cWAlCE/rocrand-7.2.4/library/include/rocrand/rocrand_normal.h Source File
rocrand_normal.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_NORMAL_H_
22#define ROCRAND_NORMAL_H_
23
28
29#include "rocrand/rocrand_lfsr113.h"
30#include "rocrand/rocrand_mrg31k3p.h"
31#include "rocrand/rocrand_mrg32k3a.h"
32#include "rocrand/rocrand_mtgp32.h"
33#include "rocrand/rocrand_philox4x32_10.h"
34#include "rocrand/rocrand_scrambled_sobol32.h"
35#include "rocrand/rocrand_scrambled_sobol64.h"
36#include "rocrand/rocrand_sobol32.h"
37#include "rocrand/rocrand_sobol64.h"
38#include "rocrand/rocrand_threefry2x32_20.h"
39#include "rocrand/rocrand_threefry2x64_20.h"
40#include "rocrand/rocrand_threefry4x32_20.h"
41#include "rocrand/rocrand_threefry4x64_20.h"
42#include "rocrand/rocrand_xorwow.h"
43
44#include "rocrand/rocrand_uniform.h"
45
46#include <hip/hip_runtime.h>
47
48#include <math.h>
49
50 namespace rocrand_device {
51 namespace detail {
52
53__forceinline__ __device__ __host__ float2 box_muller(unsigned int x, unsigned int y)
54{
55 float2 result;
56 float u = ROCRAND_2POW32_INV + (x * ROCRAND_2POW32_INV);
57 float v = ROCRAND_2POW32_INV_2PI + (y * ROCRAND_2POW32_INV_2PI);
58 float s = sqrtf(-2.0f * logf(u));
59 #ifdef __HIP_DEVICE_COMPILE__
60 __sincosf(v, &result.x, &result.y);
61 result.x *= s;
62 result.y *= s;
63 #else
64 result.x = sinf(v) * s;
65 result.y = cosf(v) * s;
66 #endif
67 return result;
68}
69
70__forceinline__ __device__ __host__ float2 box_muller(unsigned long long v)
71{
72 unsigned int x = static_cast<unsigned int>(v);
73 unsigned int y = static_cast<unsigned int>(v >> 32);
74
75 return box_muller(x, y);
76}
77
78__forceinline__ __device__ __host__ double2 box_muller_double(uint4 v)
79{
80 double2 result;
81 unsigned long long int v1 = (unsigned long long int)v.x ^
82 ((unsigned long long int)v.y << (53 - 32));
83 double u = ROCRAND_2POW53_INV_DOUBLE + (v1 * ROCRAND_2POW53_INV_DOUBLE);
84 unsigned long long int v2 = (unsigned long long int)v.z ^
85 ((unsigned long long int)v.w << (53 - 32));
86 double w = (ROCRAND_2POW53_INV_DOUBLE * 2.0) +
87 (v2 * (ROCRAND_2POW53_INV_DOUBLE * 2.0));
88 double s = sqrt(-2.0 * log(u));
89 #ifdef __HIP_DEVICE_COMPILE__
90 sincospi(w, &result.x, &result.y);
91 result.x *= s;
92 result.y *= s;
93 #else
94 result.x = sin(w * ROCRAND_PI_DOUBLE) * s;
95 result.y = cos(w * ROCRAND_PI_DOUBLE) * s;
96 #endif
97 return result;
98}
99
100__forceinline__ __device__ __host__ double2 box_muller_double(ulonglong2 v)
101{
102 unsigned int x = static_cast<unsigned int>(v.x);
103 unsigned int y = static_cast<unsigned int>(v.x >> 32);
104 unsigned int z = static_cast<unsigned int>(v.y);
105 unsigned int w = static_cast<unsigned int>(v.y >> 32);
106
107 return box_muller_double(make_uint4(x, y, z, w));
108}
109
110__forceinline__ __device__ __host__ __half2 box_muller_half(unsigned short x, unsigned short y)
111{
112 #if defined(ROCRAND_HALF_MATH_SUPPORTED)
113 __half u = __float2half(ROCRAND_2POW16_INV + (x * ROCRAND_2POW16_INV));
114 __half v = __float2half(ROCRAND_2POW16_INV_2PI + (y * ROCRAND_2POW16_INV_2PI));
115 __half s = hsqrt(__hmul(__float2half(-2.0f), hlog(u)));
116 return __half2 {
117 __hmul(hsin(v), s),
118 __hmul(hcos(v), s)
119 };
120 #else
121 float2 r;
122 float u = ROCRAND_2POW16_INV + (x * ROCRAND_2POW16_INV);
123 float v = ROCRAND_2POW16_INV_2PI + (y * ROCRAND_2POW16_INV_2PI);
124 float s = sqrtf(-2.0f * logf(u));
125 #ifdef __HIP_DEVICE_COMPILE__
126 __sincosf(v, &r.x, &r.y);
127 r.x *= s;
128 r.y *= s;
129 #else
130 r.x = sinf(v) * s;
131 r.y = cosf(v) * s;
132 #endif
133 return __half2 {
134 __float2half(r.x),
135 __float2half(r.y)
136 };
137 #endif
138}
139
140 template<typename state_type>
141__forceinline__ __device__ __host__ float2 mrg_box_muller(unsigned int x, unsigned int y)
142{
143 float2 result;
144 float u = rocrand_device::detail::mrg_uniform_distribution<state_type>(x);
145 float v = rocrand_device::detail::mrg_uniform_distribution<state_type>(y) * ROCRAND_2PI;
146 float s = sqrtf(-2.0f * logf(u));
147 #ifdef __HIP_DEVICE_COMPILE__
148 __sincosf(v, &result.x, &result.y);
149 result.x *= s;
150 result.y *= s;
151 #else
152 result.x = sinf(v) * s;
153 result.y = cosf(v) * s;
154 #endif
155 return result;
156}
157
158 template<typename state_type>
159__forceinline__ __device__ __host__ double2 mrg_box_muller_double(unsigned int x, unsigned int y)
160{
161 double2 result;
162 double u = rocrand_device::detail::mrg_uniform_distribution<state_type>(x);
163 double v = rocrand_device::detail::mrg_uniform_distribution<state_type>(y) * 2.0;
164 double s = sqrt(-2.0 * log(u));
165 #ifdef __HIP_DEVICE_COMPILE__
166 sincospi(v, &result.x, &result.y);
167 result.x *= s;
168 result.y *= s;
169 #else
170 result.x = sin(v * ROCRAND_PI_DOUBLE) * s;
171 result.y = cos(v * ROCRAND_PI_DOUBLE) * s;
172 #endif
173 return result;
174}
175
176__forceinline__ __device__ __host__ float roc_f_erfinv(float x)
177{
178 float tt1, tt2, lnx, sgn;
179 sgn = (x < 0.0f) ? -1.0f : 1.0f;
180
181 x = (1.0f - x) * (1.0f + x);
182 lnx = logf(x);
183
184 #ifdef __HIP_DEVICE_COMPILE__
185 if (isnan(lnx))
186 #else
187 if (std::isnan(lnx))
188 #endif
189 return 1.0f;
190 #ifdef __HIP_DEVICE_COMPILE__
191 else if (isinf(lnx))
192 #else
193 else if (std::isinf(lnx))
194 #endif
195 return 0.0f;
196
197 tt1 = 2.0f / (ROCRAND_PI * 0.147f) + 0.5f * lnx;
198 tt2 = 1.0f / (0.147f) * lnx;
199
200 return(sgn * sqrtf(-tt1 + sqrtf(tt1 * tt1 - tt2)));
201}
202
203__forceinline__ __device__ __host__ double roc_d_erfinv(double x)
204{
205 double tt1, tt2, lnx, sgn;
206 sgn = (x < 0.0) ? -1.0 : 1.0;
207
208 x = (1.0 - x) * (1.0 + x);
209 lnx = log(x);
210
211 #ifdef __HIP_DEVICE_COMPILE__
212 if (isnan(lnx))
213 #else
214 if (std::isnan(lnx))
215 #endif
216 return 1.0;
217 #ifdef __HIP_DEVICE_COMPILE__
218 else if (isinf(lnx))
219 #else
220 else if (std::isinf(lnx))
221 #endif
222 return 0.0;
223
224 tt1 = 2.0 / (ROCRAND_PI_DOUBLE * 0.147) + 0.5 * lnx;
225 tt2 = 1.0 / (0.147) * lnx;
226
227 return(sgn * sqrt(-tt1 + sqrt(tt1 * tt1 - tt2)));
228}
229
230__forceinline__ __device__ __host__ float normal_distribution(unsigned int x)
231{
232 float p = ::rocrand_device::detail::uniform_distribution(x);
233 float v = ROCRAND_SQRT2 * ::rocrand_device::detail::roc_f_erfinv(2.0f * p - 1.0f);
234 return v;
235}
236
237__forceinline__ __device__ __host__ float normal_distribution(unsigned long long int x)
238{
239 float p = ::rocrand_device::detail::uniform_distribution(x);
240 float v = ROCRAND_SQRT2 * ::rocrand_device::detail::roc_f_erfinv(2.0f * p - 1.0f);
241 return v;
242}
243
244__forceinline__ __device__ __host__ float2 normal_distribution2(unsigned int v1, unsigned int v2)
245{
246 return ::rocrand_device::detail::box_muller(v1, v2);
247}
248
249__forceinline__ __device__ __host__ float2 normal_distribution2(uint2 v)
250{
251 return ::rocrand_device::detail::box_muller(v.x, v.y);
252}
253
254__forceinline__ __device__ __host__ float2 normal_distribution2(unsigned long long v)
255{
256 return ::rocrand_device::detail::box_muller(v);
257}
258
259__forceinline__ __device__ __host__ float4 normal_distribution4(uint4 v)
260{
261 float2 r1 = ::rocrand_device::detail::box_muller(v.x, v.y);
262 float2 r2 = ::rocrand_device::detail::box_muller(v.z, v.w);
263 return float4{
264 r1.x,
265 r1.y,
266 r2.x,
267 r2.y
268 };
269}
270
271__forceinline__ __device__ __host__ float4 normal_distribution4(longlong2 v)
272{
273 float2 r1 = ::rocrand_device::detail::box_muller(v.x);
274 float2 r2 = ::rocrand_device::detail::box_muller(v.y);
275 return float4{r1.x, r1.y, r2.x, r2.y};
276}
277
278__forceinline__ __device__ __host__ float4 normal_distribution4(unsigned long long v1,
279 unsigned long long v2)
280{
281 float2 r1 = ::rocrand_device::detail::box_muller(v1);
282 float2 r2 = ::rocrand_device::detail::box_muller(v2);
283 return float4{r1.x, r1.y, r2.x, r2.y};
284}
285
286__forceinline__ __device__ __host__ double normal_distribution_double(unsigned int x)
287{
288 double p = ::rocrand_device::detail::uniform_distribution_double(x);
289 double v = ROCRAND_SQRT2 * ::rocrand_device::detail::roc_d_erfinv(2.0 * p - 1.0);
290 return v;
291}
292
293__forceinline__ __device__ __host__ double normal_distribution_double(unsigned long long int x)
294{
295 double p = ::rocrand_device::detail::uniform_distribution_double(x);
296 double v = ROCRAND_SQRT2 * ::rocrand_device::detail::roc_d_erfinv(2.0 * p - 1.0);
297 return v;
298}
299
300__forceinline__ __device__ __host__ double2 normal_distribution_double2(uint4 v)
301{
302 return ::rocrand_device::detail::box_muller_double(v);
303}
304
305__forceinline__ __device__ __host__ double2 normal_distribution_double2(ulonglong2 v)
306{
307 return ::rocrand_device::detail::box_muller_double(v);
308}
309
310__forceinline__ __device__ __host__ __half2 normal_distribution_half2(unsigned int v)
311{
312 return ::rocrand_device::detail::box_muller_half(
313 static_cast<unsigned short>(v),
314 static_cast<unsigned short>(v >> 16)
315 );
316}
317
318__forceinline__ __device__ __host__ __half2 normal_distribution_half2(unsigned long long v)
319{
320 return ::rocrand_device::detail::box_muller_half(static_cast<unsigned short>(v),
321 static_cast<unsigned short>(v >> 32));
322}
323
324 template<typename state_type>
325__forceinline__ __device__ __host__ float2 mrg_normal_distribution2(unsigned int v1,
326 unsigned int v2)
327{
328 return ::rocrand_device::detail::mrg_box_muller<state_type>(v1, v2);
329}
330
331 template<typename state_type>
332__forceinline__ __device__ __host__ double2 mrg_normal_distribution_double2(unsigned int v1,
333 unsigned int v2)
334{
335 return ::rocrand_device::detail::mrg_box_muller_double<state_type>(v1, v2);
336}
337
338 template<typename state_type>
339__forceinline__ __device__ __host__ __half2 mrg_normal_distribution_half2(unsigned int v)
340{
341 v = rocrand_device::detail::mrg_uniform_distribution_uint<state_type>(v);
342 return ::rocrand_device::detail::box_muller_half(
343 static_cast<unsigned short>(v),
344 static_cast<unsigned short>(v >> 16)
345 );
346}
347
348} // end namespace detail
349} // end namespace rocrand_device
350
365#ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE
366__forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_philox4x32_10* state)
367{
368 typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_philox4x32_10> bm_helper;
369
370 if(bm_helper::has_float(state))
371 {
372 return bm_helper::get_float(state);
373 }
374
375 auto state1 = rocrand(state);
376 auto state2 = rocrand(state);
377
378 float2 r = rocrand_device::detail::normal_distribution2(state1, state2);
379 bm_helper::save_float(state, r.y);
380 return r.x;
381}
382#endif // ROCRAND_DETAIL_BM_NOT_IN_STATE
383
398__forceinline__ __device__ __host__
399float2 rocrand_normal2(rocrand_state_philox4x32_10* state)
400{
401 auto state1 = rocrand(state);
402 auto state2 = rocrand(state);
403
404 return rocrand_device::detail::normal_distribution2(state1, state2);
405}
406
421__forceinline__ __device__ __host__
422float4 rocrand_normal4(rocrand_state_philox4x32_10* state)
423{
424 return rocrand_device::detail::normal_distribution4(rocrand4(state));
425}
426
441#ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE
442__forceinline__ __device__ __host__ double rocrand_normal_double(rocrand_state_philox4x32_10* state)
443{
444 typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_philox4x32_10> bm_helper;
445
446 if(bm_helper::has_double(state))
447 {
448 return bm_helper::get_double(state);
449 }
450 double2 r = rocrand_device::detail::normal_distribution_double2(rocrand4(state));
451 bm_helper::save_double(state, r.y);
452 return r.x;
453}
454#endif // ROCRAND_DETAIL_BM_NOT_IN_STATE
455
470__forceinline__ __device__ __host__
471double2 rocrand_normal_double2(rocrand_state_philox4x32_10* state)
472{
473 return rocrand_device::detail::normal_distribution_double2(rocrand4(state));
474}
475
490__forceinline__ __device__ __host__
491double4 rocrand_normal_double4(rocrand_state_philox4x32_10* state)
492{
493 double2 r1, r2;
494 r1 = rocrand_device::detail::normal_distribution_double2(rocrand4(state));
495 r2 = rocrand_device::detail::normal_distribution_double2(rocrand4(state));
496 return double4 {
497 r1.x, r1.y, r2.x, r2.y
498 };
499}
500
515#ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE
516__forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_mrg31k3p* state)
517{
518 typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_mrg31k3p> bm_helper;
519
520 if(bm_helper::has_float(state))
521 {
522 return bm_helper::get_float(state);
523 }
524
525 auto state1 = state->next();
526 auto state2 = state->next();
527
528 float2 r
529 = rocrand_device::detail::mrg_normal_distribution2<rocrand_state_mrg31k3p>(state1, state2);
530 bm_helper::save_float(state, r.y);
531 return r.x;
532}
533#endif // ROCRAND_DETAIL_BM_NOT_IN_STATE
534
549__forceinline__ __device__ __host__
550float2 rocrand_normal2(rocrand_state_mrg31k3p* state)
551{
552 auto state1 = state->next();
553 auto state2 = state->next();
554
555 return rocrand_device::detail::mrg_normal_distribution2<rocrand_state_mrg31k3p>(state1, state2);
556}
557
572#ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE
573__forceinline__ __device__ __host__ double rocrand_normal_double(rocrand_state_mrg31k3p* state)
574{
575 typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_mrg31k3p> bm_helper;
576
577 if(bm_helper::has_double(state))
578 {
579 return bm_helper::get_double(state);
580 }
581
582 auto state1 = state->next();
583 auto state2 = state->next();
584
585 double2 r
586 = rocrand_device::detail::mrg_normal_distribution_double2<rocrand_state_mrg31k3p>(state1,
587 state2);
588 bm_helper::save_double(state, r.y);
589 return r.x;
590}
591#endif // ROCRAND_DETAIL_BM_NOT_IN_STATE
592
607__forceinline__ __device__ __host__
608double2 rocrand_normal_double2(rocrand_state_mrg31k3p* state)
609{
610 auto state1 = state->next();
611 auto state2 = state->next();
612
613 return rocrand_device::detail::mrg_normal_distribution_double2<rocrand_state_mrg31k3p>(state1,
614 state2);
615}
616
631#ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE
632__forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_mrg32k3a* state)
633{
634 typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_mrg32k3a> bm_helper;
635
636 if(bm_helper::has_float(state))
637 {
638 return bm_helper::get_float(state);
639 }
640
641 auto state1 = state->next();
642 auto state2 = state->next();
643
644 float2 r
645 = rocrand_device::detail::mrg_normal_distribution2<rocrand_state_mrg32k3a>(state1, state2);
646 bm_helper::save_float(state, r.y);
647 return r.x;
648}
649#endif // ROCRAND_DETAIL_BM_NOT_IN_STATE
650
665__forceinline__ __device__ __host__
666float2 rocrand_normal2(rocrand_state_mrg32k3a* state)
667{
668 auto state1 = state->next();
669 auto state2 = state->next();
670
671 return rocrand_device::detail::mrg_normal_distribution2<rocrand_state_mrg32k3a>(state1, state2);
672}
673
688#ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE
689__forceinline__ __device__ __host__ double rocrand_normal_double(rocrand_state_mrg32k3a* state)
690{
691 typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_mrg32k3a> bm_helper;
692
693 if(bm_helper::has_double(state))
694 {
695 return bm_helper::get_double(state);
696 }
697
698 auto state1 = state->next();
699 auto state2 = state->next();
700
701 double2 r
702 = rocrand_device::detail::mrg_normal_distribution_double2<rocrand_state_mrg32k3a>(state1,
703 state2);
704 bm_helper::save_double(state, r.y);
705 return r.x;
706}
707#endif // ROCRAND_DETAIL_BM_NOT_IN_STATE
708
723__forceinline__ __device__ __host__
724double2 rocrand_normal_double2(rocrand_state_mrg32k3a* state)
725{
726 auto state1 = state->next();
727 auto state2 = state->next();
728
729 return rocrand_device::detail::mrg_normal_distribution_double2<rocrand_state_mrg32k3a>(state1,
730 state2);
731}
732
747#ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE
748__forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_xorwow* state)
749{
750 typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_xorwow> bm_helper;
751
752 if(bm_helper::has_float(state))
753 {
754 return bm_helper::get_float(state);
755 }
756 auto state1 = rocrand(state);
757 auto state2 = rocrand(state);
758 float2 r = rocrand_device::detail::normal_distribution2(state1, state2);
759 bm_helper::save_float(state, r.y);
760 return r.x;
761}
762#endif // ROCRAND_DETAIL_BM_NOT_IN_STATE
763
778__forceinline__ __device__ __host__
779float2 rocrand_normal2(rocrand_state_xorwow* state)
780{
781 auto state1 = rocrand(state);
782 auto state2 = rocrand(state);
783 return rocrand_device::detail::normal_distribution2(state1, state2);
784}
785
800#ifndef ROCRAND_DETAIL_BM_NOT_IN_STATE
801__forceinline__ __device__ __host__ double rocrand_normal_double(rocrand_state_xorwow* state)
802{
803 typedef rocrand_device::detail::engine_boxmuller_helper<rocrand_state_xorwow> bm_helper;
804
805 if(bm_helper::has_double(state))
806 {
807 return bm_helper::get_double(state);
808 }
809
810 auto state1 = rocrand(state);
811 auto state2 = rocrand(state);
812 auto state3 = rocrand(state);
813 auto state4 = rocrand(state);
814
815 double2 r = rocrand_device::detail::normal_distribution_double2(
816 uint4 { state1, state2, state3, state4 }
817 );
818 bm_helper::save_double(state, r.y);
819 return r.x;
820}
821#endif // ROCRAND_DETAIL_BM_NOT_IN_STATE
822
837__forceinline__ __device__ __host__
838double2 rocrand_normal_double2(rocrand_state_xorwow* state)
839{
840 auto state1 = rocrand(state);
841 auto state2 = rocrand(state);
842 auto state3 = rocrand(state);
843 auto state4 = rocrand(state);
844
845 return rocrand_device::detail::normal_distribution_double2(
846 uint4 { state1, state2, state3, state4 }
847 );
848}
849
862__forceinline__ __device__
863 float rocrand_normal(rocrand_state_mtgp32* state)
864{
865 return rocrand_device::detail::normal_distribution(rocrand(state));
866}
867
882__forceinline__ __device__
883float2 rocrand_normal2(rocrand_state_mtgp32* state)
884{
885 auto state1 = rocrand(state);
886 auto state2 = rocrand(state);
887 return rocrand_device::detail::normal_distribution2(state1, state2);
888}
889
902__forceinline__ __device__
903 double rocrand_normal_double(rocrand_state_mtgp32* state)
904{
905 return rocrand_device::detail::normal_distribution_double(rocrand(state));
906}
907
922__forceinline__ __device__
923double2 rocrand_normal_double2(rocrand_state_mtgp32* state)
924{
925 auto state1 = rocrand(state);
926 auto state2 = rocrand(state);
927 auto state3 = rocrand(state);
928 auto state4 = rocrand(state);
929
930 return rocrand_device::detail::normal_distribution_double2(
931 uint4{state1, state2, state3, state4});
932}
933
946__forceinline__ __device__ __host__
947 float rocrand_normal(rocrand_state_sobol32* state)
948{
949 return rocrand_device::detail::normal_distribution(rocrand(state));
950}
951
964__forceinline__ __device__ __host__
965 double rocrand_normal_double(rocrand_state_sobol32* state)
966{
967 return rocrand_device::detail::normal_distribution_double(rocrand(state));
968}
969
982__forceinline__ __device__ __host__
983 float rocrand_normal(rocrand_state_scrambled_sobol32* state)
984{
985 return rocrand_device::detail::normal_distribution(rocrand(state));
986}
987
1000__forceinline__ __device__ __host__
1001 double rocrand_normal_double(rocrand_state_scrambled_sobol32* state)
1002{
1003 return rocrand_device::detail::normal_distribution_double(rocrand(state));
1004}
1005
1018__forceinline__ __device__ __host__
1019 float rocrand_normal(rocrand_state_sobol64* state)
1020{
1021 return rocrand_device::detail::normal_distribution(rocrand(state));
1022}
1023
1036__forceinline__ __device__ __host__
1037 double rocrand_normal_double(rocrand_state_sobol64* state)
1038{
1039 return rocrand_device::detail::normal_distribution_double(rocrand(state));
1040}
1041
1054__forceinline__ __device__ __host__
1055 float rocrand_normal(rocrand_state_scrambled_sobol64* state)
1056{
1057 return rocrand_device::detail::normal_distribution(rocrand(state));
1058}
1059
1072__forceinline__ __device__ __host__
1073 double rocrand_normal_double(rocrand_state_scrambled_sobol64* state)
1074{
1075 return rocrand_device::detail::normal_distribution_double(rocrand(state));
1076}
1077
1090__forceinline__ __device__ __host__
1091 float rocrand_normal(rocrand_state_lfsr113* state)
1092{
1093 return rocrand_device::detail::normal_distribution(rocrand(state));
1094}
1095
1110__forceinline__ __device__ __host__
1111float2 rocrand_normal2(rocrand_state_lfsr113* state)
1112{
1113 auto state1 = rocrand(state);
1114 auto state2 = rocrand(state);
1115
1116 return rocrand_device::detail::normal_distribution2(state1, state2);
1117}
1118
1131__forceinline__ __device__ __host__
1132 double rocrand_normal_double(rocrand_state_lfsr113* state)
1133{
1134 return rocrand_device::detail::normal_distribution_double(rocrand(state));
1135}
1136
1151__forceinline__ __device__ __host__
1152double2 rocrand_normal_double2(rocrand_state_lfsr113* state)
1153{
1154 auto state1 = rocrand(state);
1155 auto state2 = rocrand(state);
1156 auto state3 = rocrand(state);
1157 auto state4 = rocrand(state);
1158
1159 return rocrand_device::detail::normal_distribution_double2(
1160 uint4{state1, state2, state3, state4});
1161}
1162
1175__forceinline__ __device__ __host__
1176 float rocrand_normal(rocrand_state_threefry2x32_20* state)
1177{
1178 return rocrand_device::detail::normal_distribution(rocrand(state));
1179}
1180
1195__forceinline__ __device__ __host__
1196float2 rocrand_normal2(rocrand_state_threefry2x32_20* state)
1197{
1198 return rocrand_device::detail::normal_distribution2(rocrand2(state));
1199}
1200
1213__forceinline__ __device__ __host__
1214 double rocrand_normal_double(rocrand_state_threefry2x32_20* state)
1215{
1216 return rocrand_device::detail::normal_distribution_double(rocrand(state));
1217}
1218
1233__forceinline__ __device__ __host__
1234double2 rocrand_normal_double2(rocrand_state_threefry2x32_20* state)
1235{
1236 auto state1 = rocrand2(state);
1237 auto state2 = rocrand2(state);
1238
1239 return rocrand_device::detail::normal_distribution_double2(
1240 uint4{state1.x, state1.y, state2.x, state2.y});
1241}
1242
1255__forceinline__ __device__ __host__
1256 float rocrand_normal(rocrand_state_threefry2x64_20* state)
1257{
1258 return rocrand_device::detail::normal_distribution(rocrand(state));
1259}
1260
1275__forceinline__ __device__ __host__
1276float2 rocrand_normal2(rocrand_state_threefry2x64_20* state)
1277{
1278 return rocrand_device::detail::normal_distribution2(rocrand(state));
1279}
1280
1293__forceinline__ __device__ __host__
1294 double rocrand_normal_double(rocrand_state_threefry2x64_20* state)
1295{
1296 return rocrand_device::detail::normal_distribution_double(rocrand(state));
1297}
1298
1313__forceinline__ __device__ __host__
1314double2 rocrand_normal_double2(rocrand_state_threefry2x64_20* state)
1315{
1316 return rocrand_device::detail::normal_distribution_double2(rocrand2(state));
1317}
1318
1331__forceinline__ __device__ __host__
1332 float rocrand_normal(rocrand_state_threefry4x32_20* state)
1333{
1334 return rocrand_device::detail::normal_distribution(rocrand(state));
1335}
1336
1351__forceinline__ __device__ __host__
1352float2 rocrand_normal2(rocrand_state_threefry4x32_20* state)
1353{
1354 auto state1 = rocrand(state);
1355 auto state2 = rocrand(state);
1356
1357 return rocrand_device::detail::normal_distribution2(state1, state2);
1358}
1359
1372__forceinline__ __device__ __host__
1373 double rocrand_normal_double(rocrand_state_threefry4x32_20* state)
1374{
1375 return rocrand_device::detail::normal_distribution_double(rocrand(state));
1376}
1377
1392__forceinline__ __device__ __host__
1393double2 rocrand_normal_double2(rocrand_state_threefry4x32_20* state)
1394{
1395 return rocrand_device::detail::normal_distribution_double2(rocrand4(state));
1396}
1397
1410__forceinline__ __device__ __host__
1411 float rocrand_normal(rocrand_state_threefry4x64_20* state)
1412{
1413 return rocrand_device::detail::normal_distribution(rocrand(state));
1414}
1415
1430__forceinline__ __device__ __host__
1431float2 rocrand_normal2(rocrand_state_threefry4x64_20* state)
1432{
1433 auto state1 = rocrand(state);
1434 auto state2 = rocrand(state);
1435
1436 return rocrand_device::detail::normal_distribution2(state1, state2);
1437}
1438
1451__forceinline__ __device__ __host__
1452 double rocrand_normal_double(rocrand_state_threefry4x64_20* state)
1453{
1454 return rocrand_device::detail::normal_distribution_double(rocrand(state));
1455}
1456
1471__forceinline__ __device__ __host__
1472double2 rocrand_normal_double2(rocrand_state_threefry4x64_20* state)
1473{
1474 auto state1 = rocrand(state);
1475 auto state2 = rocrand(state);
1476
1477 return rocrand_device::detail::normal_distribution_double2(ulonglong2{state1, state2});
1478}
1479 // end of group rocranddevice
1481
1482#endif // ROCRAND_NORMAL_H_
rocrand_normal_double4
__forceinline__ __device__ __host__ double4 rocrand_normal_double4(rocrand_state_philox4x32_10 *state)
Returns four normally distributed double values.
Definition rocrand_normal.h:491
rocrand_normal_double2
__forceinline__ __device__ __host__ double2 rocrand_normal_double2(rocrand_state_philox4x32_10 *state)
Returns two normally distributed double values.
Definition rocrand_normal.h:471
rocrand_normal
__forceinline__ __device__ __host__ float rocrand_normal(rocrand_state_philox4x32_10 *state)
Returns a normally distributed float value.
Definition rocrand_normal.h:366
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_normal_double
__forceinline__ __device__ __host__ double rocrand_normal_double(rocrand_state_philox4x32_10 *state)
Returns a normally distributed double value.
Definition rocrand_normal.h:442
rocrand_normal4
__forceinline__ __device__ __host__ float4 rocrand_normal4(rocrand_state_philox4x32_10 *state)
Returns four normally distributed float values.
Definition rocrand_normal.h:422
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_normal2
__forceinline__ __device__ __host__ float2 rocrand_normal2(rocrand_state_philox4x32_10 *state)
Returns two normally distributed float values.
Definition rocrand_normal.h:399