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

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

API library: /build/reproducible-path/rocrand-6.4.1/library/include/rocrand/rocrand_uniform.h Source File
rocrand_uniform.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
26#ifndef ROCRAND_UNIFORM_H_
27#define ROCRAND_UNIFORM_H_
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_common.h"
45
46namespace rocrand_device {
47namespace detail {
48
49struct two_uints
50{
51 unsigned int x;
52 unsigned int y;
53};
54
55union two_uints_to_ulong
56{
57 two_uints uint2_value;
58 unsigned long long int ulong_value;
59};
60
61// For unsigned integer between 0 and UINT_MAX, returns value between
62// 0.0f and 1.0f, excluding 0.0f and including 1.0f.
63__forceinline__ __device__ __host__ float uniform_distribution(unsigned int v)
64{
65 return ROCRAND_2POW32_INV + (v * ROCRAND_2POW32_INV);
66}
67
68// For unsigned integer between 0 and ULLONG_MAX, returns value between
69// 0.0f and 1.0f, excluding 0.0f and including 1.0f.
70__forceinline__ __device__ __host__ float uniform_distribution(unsigned long long int v)
71{
72 return ROCRAND_2POW32_INV + (v >> 32) * ROCRAND_2POW32_INV;
73}
74
75__forceinline__ __device__ __host__ float4 uniform_distribution4(uint4 v)
76{
77 return float4 {
78 ROCRAND_2POW32_INV + (v.x * ROCRAND_2POW32_INV),
79 ROCRAND_2POW32_INV + (v.y * ROCRAND_2POW32_INV),
80 ROCRAND_2POW32_INV + (v.z * ROCRAND_2POW32_INV),
81 ROCRAND_2POW32_INV + (v.w * ROCRAND_2POW32_INV)
82 };
83}
84
85__forceinline__ __device__ __host__ float4 uniform_distribution4(ulonglong4 v)
86{
87 return float4{ROCRAND_2POW64_INV + (v.x * ROCRAND_2POW64_INV),
88 ROCRAND_2POW64_INV + (v.y * ROCRAND_2POW64_INV),
89 ROCRAND_2POW64_INV + (v.z * ROCRAND_2POW64_INV),
90 ROCRAND_2POW64_INV + (v.w * ROCRAND_2POW64_INV)};
91}
92
93// For unsigned integer between 0 and UINT_MAX, returns value between
94// 0.0 and 1.0, excluding 0.0 and including 1.0.
95__forceinline__ __device__ __host__ double uniform_distribution_double(unsigned int v)
96{
97 return ROCRAND_2POW32_INV_DOUBLE + (v * ROCRAND_2POW32_INV_DOUBLE);
98}
99
100__forceinline__ __device__ __host__ double uniform_distribution_double(unsigned int v1,
101 unsigned int v2)
102{
103 two_uints_to_ulong v;
104 v.uint2_value.x = v1;
105 v.uint2_value.y = (v2 >> 11);
106 return ROCRAND_2POW53_INV_DOUBLE + (v.ulong_value * ROCRAND_2POW53_INV_DOUBLE);
107}
108
109__forceinline__ __device__ __host__ double uniform_distribution_double(unsigned long long int v)
110{
111 return ROCRAND_2POW53_INV_DOUBLE + (
112 // 2^53 is the biggest int that can be stored in double, such
113 // that it and all smaller integers can be stored in double
114 (v >> 11) * ROCRAND_2POW53_INV_DOUBLE
115 );
116}
117
118__forceinline__ __device__ __host__ double2 uniform_distribution_double2(uint4 v)
119{
120 return double2 {
121 uniform_distribution_double(v.x, v.y),
122 uniform_distribution_double(v.z, v.w)
123 };
124}
125
126__forceinline__ __device__ __host__ double4 uniform_distribution_double4(uint4 v1, uint4 v2)
127{
128 return double4 {
129 uniform_distribution_double(v1.x, v1.y),
130 uniform_distribution_double(v1.z, v1.w),
131 uniform_distribution_double(v2.x, v2.y),
132 uniform_distribution_double(v2.z, v2.w)
133 };
134}
135
136__forceinline__ __device__ __host__ double2 uniform_distribution_double2(ulonglong2 v)
137{
138 return double2{uniform_distribution_double(v.x), uniform_distribution_double(v.y)};
139}
140
141__forceinline__ __device__ __host__ double2 uniform_distribution_double2(ulonglong4 v)
142{
143 return double2{uniform_distribution_double(v.x), uniform_distribution_double(v.y)};
144}
145
146__forceinline__ __device__ __host__ double4 uniform_distribution_double4(ulonglong4 v)
147{
148 return double4{uniform_distribution_double(v.x),
149 uniform_distribution_double(v.z),
150 uniform_distribution_double(v.x),
151 uniform_distribution_double(v.z)};
152}
153
154__forceinline__ __device__ __host__ __half uniform_distribution_half(unsigned short v)
155{
156 return __float2half(ROCRAND_2POW16_INV + (v * ROCRAND_2POW16_INV));
157}
158
159// For an unsigned integer produced by an MRG-based engine, returns a value
160// in range [0, UINT32_MAX].
161template<typename state_type>
162__forceinline__ __device__ __host__ unsigned int mrg_uniform_distribution_uint(unsigned int v)
163 = delete;
164
165template<>
166__forceinline__ __device__ __host__ unsigned int
167 mrg_uniform_distribution_uint<rocrand_state_mrg31k3p>(unsigned int v)
168{
169 return static_cast<unsigned int>((v - 1) * ROCRAND_MRG31K3P_UINT32_NORM);
170}
171
172template<>
173__forceinline__ __device__ __host__ unsigned int
174 mrg_uniform_distribution_uint<rocrand_state_mrg32k3a>(unsigned int v)
175{
176 return static_cast<unsigned int>((v - 1) * ROCRAND_MRG32K3A_UINT_NORM);
177}
178
179// For an unsigned integer produced by an MRG-based engine, returns value between
180// 0.0f and 1.0f, excluding 0.0f and including 1.0f.
181template<typename state_type>
182__forceinline__ __device__ __host__ float mrg_uniform_distribution(unsigned int v) = delete;
183
184template<>
185__forceinline__ __device__ __host__ float
186 mrg_uniform_distribution<rocrand_state_mrg31k3p>(unsigned int v)
187{
188 double ret = static_cast<double>(v) * ROCRAND_MRG31K3P_NORM_DOUBLE;
189 return static_cast<float>(ret);
190}
191
192template<>
193__forceinline__ __device__ __host__ float
194 mrg_uniform_distribution<rocrand_state_mrg32k3a>(unsigned int v)
195{
196 double ret = static_cast<double>(v) * ROCRAND_MRG32K3A_NORM_DOUBLE;
197 return static_cast<float>(ret);
198}
199
200// For an unsigned integer produced by an MRG generator, returns value between
201// 0.0 and 1.0, excluding 0.0 and including 1.0.
202template<typename state_type>
203__forceinline__ __device__ __host__ double mrg_uniform_distribution_double(unsigned int v) = delete;
204
205template<>
206__forceinline__ __device__ __host__ double
207 mrg_uniform_distribution_double<rocrand_state_mrg31k3p>(unsigned int v)
208{
209 double ret = static_cast<double>(v) * ROCRAND_MRG31K3P_NORM_DOUBLE;
210 return ret;
211}
212
213template<>
214__forceinline__ __device__ __host__ double
215 mrg_uniform_distribution_double<rocrand_state_mrg32k3a>(unsigned int v)
216{
217 double ret = static_cast<double>(v) * ROCRAND_MRG32K3A_NORM_DOUBLE;
218 return ret;
219}
220
221} // end namespace detail
222} // end namespace rocrand_device
223
236__forceinline__ __device__ __host__ float rocrand_uniform(rocrand_state_philox4x32_10* state)
237{
238 return rocrand_device::detail::uniform_distribution(rocrand(state));
239}
240
253__forceinline__ __device__ __host__ float2 rocrand_uniform2(rocrand_state_philox4x32_10* state)
254{
255 auto state1 = rocrand(state);
256 auto state2 = rocrand(state);
257
258 return float2 {
259 rocrand_device::detail::uniform_distribution(state1),
260 rocrand_device::detail::uniform_distribution(state2)
261 };
262}
263
276__forceinline__ __device__ __host__ float4 rocrand_uniform4(rocrand_state_philox4x32_10* state)
277{
278 return rocrand_device::detail::uniform_distribution4(rocrand4(state));
279}
280
293__forceinline__ __device__ __host__ double
294 rocrand_uniform_double(rocrand_state_philox4x32_10* state)
295{
296 auto state1 = rocrand(state);
297 auto state2 = rocrand(state);
298
299 return rocrand_device::detail::uniform_distribution_double(state1, state2);
300}
301
314__forceinline__ __device__ __host__ double2
315 rocrand_uniform_double2(rocrand_state_philox4x32_10* state)
316{
317 return rocrand_device::detail::uniform_distribution_double2(rocrand4(state));
318}
319
332__forceinline__ __device__ __host__ double4
333 rocrand_uniform_double4(rocrand_state_philox4x32_10* state)
334{
335 return rocrand_device::detail::uniform_distribution_double4(rocrand4(state), rocrand4(state));
336}
337
350__forceinline__ __device__ __host__ float rocrand_uniform(rocrand_state_mrg31k3p* state)
351{
352 return rocrand_device::detail::mrg_uniform_distribution<rocrand_state_mrg31k3p>(state->next());
353}
354
370__forceinline__ __device__ __host__ double rocrand_uniform_double(rocrand_state_mrg31k3p* state)
371{
372 return rocrand_device::detail::mrg_uniform_distribution_double<rocrand_state_mrg31k3p>(
373 state->next());
374}
375
388__forceinline__ __device__ __host__ float rocrand_uniform(rocrand_state_mrg32k3a* state)
389{
390 return rocrand_device::detail::mrg_uniform_distribution<rocrand_state_mrg32k3a>(state->next());
391}
392
408__forceinline__ __device__ __host__ double rocrand_uniform_double(rocrand_state_mrg32k3a* state)
409{
410 return rocrand_device::detail::mrg_uniform_distribution_double<rocrand_state_mrg32k3a>(
411 state->next());
412}
413
426__forceinline__ __device__ __host__ float rocrand_uniform(rocrand_state_xorwow* state)
427{
428 return rocrand_device::detail::uniform_distribution(rocrand(state));
429}
430
443__forceinline__ __device__ __host__ double rocrand_uniform_double(rocrand_state_xorwow* state)
444{
445 auto state1 = rocrand(state);
446 auto state2 = rocrand(state);
447
448 return rocrand_device::detail::uniform_distribution_double(state1, state2);
449}
450
463__forceinline__ __device__ float rocrand_uniform(rocrand_state_mtgp32* state)
464{
465 return rocrand_device::detail::uniform_distribution(rocrand(state));
466}
467
483__forceinline__ __device__ double rocrand_uniform_double(rocrand_state_mtgp32* state)
484{
485 return rocrand_device::detail::uniform_distribution_double(rocrand(state));
486}
487
500__forceinline__ __device__ __host__ float rocrand_uniform(rocrand_state_sobol32* state)
501{
502 return rocrand_device::detail::uniform_distribution(rocrand(state));
503}
504
520__forceinline__ __device__ __host__ double rocrand_uniform_double(rocrand_state_sobol32* state)
521{
522 return rocrand_device::detail::uniform_distribution_double(rocrand(state));
523}
524
537__forceinline__ __device__ __host__ float rocrand_uniform(rocrand_state_scrambled_sobol32* state)
538{
539 return rocrand_device::detail::uniform_distribution(rocrand(state));
540}
541
557__forceinline__ __device__ __host__ double
558 rocrand_uniform_double(rocrand_state_scrambled_sobol32* state)
559{
560 return rocrand_device::detail::uniform_distribution_double(rocrand(state));
561}
562
575__forceinline__ __device__ __host__ float rocrand_uniform(rocrand_state_sobol64* state)
576{
577 return rocrand_device::detail::uniform_distribution(rocrand(state));
578}
579
592__forceinline__ __device__ __host__ double rocrand_uniform_double(rocrand_state_sobol64* state)
593{
594 return rocrand_device::detail::uniform_distribution_double(rocrand(state));
595}
596
609__forceinline__ __device__ __host__ float rocrand_uniform(rocrand_state_scrambled_sobol64* state)
610{
611 return rocrand_device::detail::uniform_distribution(rocrand(state));
612}
613
626__forceinline__ __device__ __host__ double
627 rocrand_uniform_double(rocrand_state_scrambled_sobol64* state)
628{
629 return rocrand_device::detail::uniform_distribution_double(rocrand(state));
630}
631
644__forceinline__ __device__ __host__ float rocrand_uniform(rocrand_state_lfsr113* state)
645{
646 return rocrand_device::detail::uniform_distribution(rocrand(state));
647}
648
664__forceinline__ __device__ __host__ double rocrand_uniform_double(rocrand_state_lfsr113* state)
665{
666 return rocrand_device::detail::uniform_distribution_double(rocrand(state));
667}
668
681__forceinline__ __device__ __host__ float rocrand_uniform(rocrand_state_threefry2x32_20* state)
682{
683 return rocrand_device::detail::uniform_distribution(rocrand(state));
684}
685
701__forceinline__ __device__ __host__ double
702 rocrand_uniform_double(rocrand_state_threefry2x32_20* state)
703{
704 return rocrand_device::detail::uniform_distribution_double(rocrand(state));
705}
706
719__forceinline__ __device__ __host__ float rocrand_uniform(rocrand_state_threefry2x64_20* state)
720{
721 return rocrand_device::detail::uniform_distribution(rocrand(state));
722}
723
739__forceinline__ __device__ __host__ double
740 rocrand_uniform_double(rocrand_state_threefry2x64_20* state)
741{
742 return rocrand_device::detail::uniform_distribution_double(rocrand(state));
743}
744
757__forceinline__ __device__ __host__ float rocrand_uniform(rocrand_state_threefry4x32_20* state)
758{
759 return rocrand_device::detail::uniform_distribution(rocrand(state));
760}
761
777__forceinline__ __device__ __host__ double
778 rocrand_uniform_double(rocrand_state_threefry4x32_20* state)
779{
780 return rocrand_device::detail::uniform_distribution_double(rocrand(state));
781}
782
795__forceinline__ __device__ __host__ float rocrand_uniform(rocrand_state_threefry4x64_20* state)
796{
797 return rocrand_device::detail::uniform_distribution(rocrand(state));
798}
799
815__forceinline__ __device__ __host__ double
816 rocrand_uniform_double(rocrand_state_threefry4x64_20* state)
817{
818 return rocrand_device::detail::uniform_distribution_double(rocrand(state));
819}
820
// end of group rocranddevice
822
823#endif // ROCRAND_UNIFORM_H_
__forceinline__ __device__ __host__ float4 rocrand_uniform4(rocrand_state_philox4x32_10 *state)
Returns four uniformly distributed random float values from (0; 1] range.
Definition rocrand_uniform.h:276
__forceinline__ __device__ __host__ double4 rocrand_uniform_double4(rocrand_state_philox4x32_10 *state)
Returns four uniformly distributed random double values from (0; 1] range.
Definition rocrand_uniform.h:333
__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_uniform_double(rocrand_state_philox4x32_10 *state)
Returns a uniformly distributed random double value from (0; 1] range.
Definition rocrand_uniform.h:294
__forceinline__ __device__ __host__ float rocrand_uniform(rocrand_state_philox4x32_10 *state)
Returns a uniformly distributed random float value from (0; 1] range.
Definition rocrand_uniform.h:236
__forceinline__ __device__ __host__ float2 rocrand_uniform2(rocrand_state_philox4x32_10 *state)
Returns two uniformly distributed random float values from (0; 1] range.
Definition rocrand_uniform.h:253
__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__ double2 rocrand_uniform_double2(rocrand_state_philox4x32_10 *state)
Returns two uniformly distributed random double values from (0; 1] range.
Definition rocrand_uniform.h:315