Skip to main content
`Ctrl`+`K`

# /build/rocrand-7S8maf/rocrand-7.1.1/library/include/rocrand/rocrand_uniform.h Source File

API library: /build/rocrand-7S8maf/rocrand-7.1.1/library/include/rocrand/rocrand_uniform.h Source File
rocrand_uniform.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
25
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
46#include <hip/hip_runtime.h>
47
48 namespace rocrand_device {
49 namespace detail {
50
51 struct two_uints
52{
53 unsigned int x;
54 unsigned int y;
55};
56
57 union two_uints_to_ulong
58{
59 two_uints uint2_value;
60 unsigned long long int ulong_value;
61};
62
63// For unsigned integer between 0 and UINT_MAX, returns value between
64// 0.0f and 1.0f, excluding 0.0f and including 1.0f.
65__forceinline__ __device__ __host__ float uniform_distribution(unsigned int v)
66{
67 return ROCRAND_2POW32_INV + (v * ROCRAND_2POW32_INV);
68}
69
70// For unsigned integer between 0 and ULLONG_MAX, returns value between
71// 0.0f and 1.0f, excluding 0.0f and including 1.0f.
72__forceinline__ __device__ __host__ float uniform_distribution(unsigned long long int v)
73{
74 return ROCRAND_2POW32_INV + (v >> 32) * ROCRAND_2POW32_INV;
75}
76
77__forceinline__ __device__ __host__ float4 uniform_distribution4(uint4 v)
78{
79 return float4 {
80 ROCRAND_2POW32_INV + (v.x * ROCRAND_2POW32_INV),
81 ROCRAND_2POW32_INV + (v.y * ROCRAND_2POW32_INV),
82 ROCRAND_2POW32_INV + (v.z * ROCRAND_2POW32_INV),
83 ROCRAND_2POW32_INV + (v.w * ROCRAND_2POW32_INV)
84 };
85}
86
87__forceinline__ __device__ __host__ float4 uniform_distribution4(ulonglong4 v)
88{
89 return float4{ROCRAND_2POW64_INV + (v.x * ROCRAND_2POW64_INV),
90 ROCRAND_2POW64_INV + (v.y * ROCRAND_2POW64_INV),
91 ROCRAND_2POW64_INV + (v.z * ROCRAND_2POW64_INV),
92 ROCRAND_2POW64_INV + (v.w * ROCRAND_2POW64_INV)};
93}
94
95// For unsigned integer between 0 and UINT_MAX, returns value between
96// 0.0 and 1.0, excluding 0.0 and including 1.0.
97__forceinline__ __device__ __host__ double uniform_distribution_double(unsigned int v)
98{
99 return ROCRAND_2POW32_INV_DOUBLE + (v * ROCRAND_2POW32_INV_DOUBLE);
100}
101
102__forceinline__ __device__ __host__ double uniform_distribution_double(unsigned int v1,
103 unsigned int v2)
104{
105 two_uints_to_ulong v;
106 v.uint2_value.x = v1;
107 v.uint2_value.y = (v2 >> 11);
108 return ROCRAND_2POW53_INV_DOUBLE + (v.ulong_value * ROCRAND_2POW53_INV_DOUBLE);
109}
110
111__forceinline__ __device__ __host__ double uniform_distribution_double(unsigned long long int v)
112{
113 return ROCRAND_2POW53_INV_DOUBLE + (
114 // 2^53 is the biggest int that can be stored in double, such
115 // that it and all smaller integers can be stored in double
116 (v >> 11) * ROCRAND_2POW53_INV_DOUBLE
117 );
118}
119
120__forceinline__ __device__ __host__ double2 uniform_distribution_double2(uint4 v)
121{
122 return double2 {
123 uniform_distribution_double(v.x, v.y),
124 uniform_distribution_double(v.z, v.w)
125 };
126}
127
128__forceinline__ __device__ __host__ double4 uniform_distribution_double4(uint4 v1, uint4 v2)
129{
130 return double4 {
131 uniform_distribution_double(v1.x, v1.y),
132 uniform_distribution_double(v1.z, v1.w),
133 uniform_distribution_double(v2.x, v2.y),
134 uniform_distribution_double(v2.z, v2.w)
135 };
136}
137
138__forceinline__ __device__ __host__ double2 uniform_distribution_double2(ulonglong2 v)
139{
140 return double2{uniform_distribution_double(v.x), uniform_distribution_double(v.y)};
141}
142
143__forceinline__ __device__ __host__ double2 uniform_distribution_double2(ulonglong4 v)
144{
145 return double2{uniform_distribution_double(v.x), uniform_distribution_double(v.y)};
146}
147
148__forceinline__ __device__ __host__ double4 uniform_distribution_double4(ulonglong4 v)
149{
150 return double4{uniform_distribution_double(v.x),
151 uniform_distribution_double(v.z),
152 uniform_distribution_double(v.x),
153 uniform_distribution_double(v.z)};
154}
155
156__forceinline__ __device__ __host__ __half uniform_distribution_half(unsigned short v)
157{
158 return __float2half(ROCRAND_2POW16_INV + (v * ROCRAND_2POW16_INV));
159}
160
161// For an unsigned integer produced by an MRG-based engine, returns a value
162// in range [0, UINT32_MAX].
163 template<typename state_type>
164__forceinline__ __device__ __host__ unsigned int mrg_uniform_distribution_uint(unsigned int v)
165 = delete;
166
167 template<>
168__forceinline__ __device__ __host__ unsigned int
169 mrg_uniform_distribution_uint<rocrand_state_mrg31k3p>(unsigned int v)
170{
171 return static_cast<unsigned int>((v - 1) * ROCRAND_MRG31K3P_UINT32_NORM);
172}
173
174 template<>
175__forceinline__ __device__ __host__ unsigned int
176 mrg_uniform_distribution_uint<rocrand_state_mrg32k3a>(unsigned int v)
177{
178 return static_cast<unsigned int>((v - 1) * ROCRAND_MRG32K3A_UINT_NORM);
179}
180
181// For an unsigned integer produced by an MRG-based engine, returns value between
182// 0.0f and 1.0f, excluding 0.0f and including 1.0f.
183 template<typename state_type>
184__forceinline__ __device__ __host__ float mrg_uniform_distribution(unsigned int v) = delete;
185
186 template<>
187__forceinline__ __device__ __host__ float
188 mrg_uniform_distribution<rocrand_state_mrg31k3p>(unsigned int v)
189{
190 double ret = static_cast<double>(v) * ROCRAND_MRG31K3P_NORM_DOUBLE;
191 return static_cast<float>(ret);
192}
193
194 template<>
195__forceinline__ __device__ __host__ float
196 mrg_uniform_distribution<rocrand_state_mrg32k3a>(unsigned int v)
197{
198 double ret = static_cast<double>(v) * ROCRAND_MRG32K3A_NORM_DOUBLE;
199 return static_cast<float>(ret);
200}
201
202// For an unsigned integer produced by an MRG generator, returns value between
203// 0.0 and 1.0, excluding 0.0 and including 1.0.
204 template<typename state_type>
205__forceinline__ __device__ __host__ double mrg_uniform_distribution_double(unsigned int v) = delete;
206
207 template<>
208__forceinline__ __device__ __host__ double
209 mrg_uniform_distribution_double<rocrand_state_mrg31k3p>(unsigned int v)
210{
211 double ret = static_cast<double>(v) * ROCRAND_MRG31K3P_NORM_DOUBLE;
212 return ret;
213}
214
215 template<>
216__forceinline__ __device__ __host__ double
217 mrg_uniform_distribution_double<rocrand_state_mrg32k3a>(unsigned int v)
218{
219 double ret = static_cast<double>(v) * ROCRAND_MRG32K3A_NORM_DOUBLE;
220 return ret;
221}
222
223} // end namespace detail
224} // end namespace rocrand_device
225
238__forceinline__ __device__ __host__
239 float rocrand_uniform(rocrand_state_philox4x32_10* state)
240{
241 return rocrand_device::detail::uniform_distribution(rocrand(state));
242}
243
256__forceinline__ __device__ __host__
257float2 rocrand_uniform2(rocrand_state_philox4x32_10* state)
258{
259 auto state1 = rocrand(state);
260 auto state2 = rocrand(state);
261
262 return float2 {
263 rocrand_device::detail::uniform_distribution(state1),
264 rocrand_device::detail::uniform_distribution(state2)
265 };
266}
267
280__forceinline__ __device__ __host__
281float4 rocrand_uniform4(rocrand_state_philox4x32_10* state)
282{
283 return rocrand_device::detail::uniform_distribution4(rocrand4(state));
284}
285
298__forceinline__ __device__ __host__
299 double rocrand_uniform_double(rocrand_state_philox4x32_10* state)
300{
301 auto state1 = rocrand(state);
302 auto state2 = rocrand(state);
303
304 return rocrand_device::detail::uniform_distribution_double(state1, state2);
305}
306
319__forceinline__ __device__ __host__
320double2 rocrand_uniform_double2(rocrand_state_philox4x32_10* state)
321{
322 return rocrand_device::detail::uniform_distribution_double2(rocrand4(state));
323}
324
337__forceinline__ __device__ __host__
338double4 rocrand_uniform_double4(rocrand_state_philox4x32_10* state)
339{
340 return rocrand_device::detail::uniform_distribution_double4(rocrand4(state), rocrand4(state));
341}
342
355__forceinline__ __device__ __host__
356 float rocrand_uniform(rocrand_state_mrg31k3p* state)
357{
358 return rocrand_device::detail::mrg_uniform_distribution<rocrand_state_mrg31k3p>(state->next());
359}
360
376__forceinline__ __device__ __host__
377 double rocrand_uniform_double(rocrand_state_mrg31k3p* state)
378{
379 return rocrand_device::detail::mrg_uniform_distribution_double<rocrand_state_mrg31k3p>(
380 state->next());
381}
382
395__forceinline__ __device__ __host__
396 float rocrand_uniform(rocrand_state_mrg32k3a* state)
397{
398 return rocrand_device::detail::mrg_uniform_distribution<rocrand_state_mrg32k3a>(state->next());
399}
400
416__forceinline__ __device__ __host__
417 double rocrand_uniform_double(rocrand_state_mrg32k3a* state)
418{
419 return rocrand_device::detail::mrg_uniform_distribution_double<rocrand_state_mrg32k3a>(
420 state->next());
421}
422
435__forceinline__ __device__ __host__
436 float rocrand_uniform(rocrand_state_xorwow* state)
437{
438 return rocrand_device::detail::uniform_distribution(rocrand(state));
439}
440
453__forceinline__ __device__ __host__
454 double rocrand_uniform_double(rocrand_state_xorwow* state)
455{
456 auto state1 = rocrand(state);
457 auto state2 = rocrand(state);
458
459 return rocrand_device::detail::uniform_distribution_double(state1, state2);
460}
461
474__forceinline__ __device__
475 float rocrand_uniform(rocrand_state_mtgp32* state)
476{
477 return rocrand_device::detail::uniform_distribution(rocrand(state));
478}
479
495__forceinline__ __device__
496 double rocrand_uniform_double(rocrand_state_mtgp32* state)
497{
498 return rocrand_device::detail::uniform_distribution_double(rocrand(state));
499}
500
513__forceinline__ __device__ __host__
514 float rocrand_uniform(rocrand_state_sobol32* state)
515{
516 return rocrand_device::detail::uniform_distribution(rocrand(state));
517}
518
534__forceinline__ __device__ __host__
535 double rocrand_uniform_double(rocrand_state_sobol32* state)
536{
537 return rocrand_device::detail::uniform_distribution_double(rocrand(state));
538}
539
552__forceinline__ __device__ __host__
553 float rocrand_uniform(rocrand_state_scrambled_sobol32* state)
554{
555 return rocrand_device::detail::uniform_distribution(rocrand(state));
556}
557
573__forceinline__ __device__ __host__
574 double rocrand_uniform_double(rocrand_state_scrambled_sobol32* state)
575{
576 return rocrand_device::detail::uniform_distribution_double(rocrand(state));
577}
578
591__forceinline__ __device__ __host__
592 float rocrand_uniform(rocrand_state_sobol64* state)
593{
594 return rocrand_device::detail::uniform_distribution(rocrand(state));
595}
596
609__forceinline__ __device__ __host__
610 double rocrand_uniform_double(rocrand_state_sobol64* state)
611{
612 return rocrand_device::detail::uniform_distribution_double(rocrand(state));
613}
614
627__forceinline__ __device__ __host__
628 float rocrand_uniform(rocrand_state_scrambled_sobol64* state)
629{
630 return rocrand_device::detail::uniform_distribution(rocrand(state));
631}
632
645__forceinline__ __device__ __host__
646 double rocrand_uniform_double(rocrand_state_scrambled_sobol64* state)
647{
648 return rocrand_device::detail::uniform_distribution_double(rocrand(state));
649}
650
663__forceinline__ __device__ __host__
664 float rocrand_uniform(rocrand_state_lfsr113* state)
665{
666 return rocrand_device::detail::uniform_distribution(rocrand(state));
667}
668
684__forceinline__ __device__ __host__
685 double rocrand_uniform_double(rocrand_state_lfsr113* state)
686{
687 return rocrand_device::detail::uniform_distribution_double(rocrand(state));
688}
689
702__forceinline__ __device__ __host__
703 float rocrand_uniform(rocrand_state_threefry2x32_20* state)
704{
705 return rocrand_device::detail::uniform_distribution(rocrand(state));
706}
707
723__forceinline__ __device__ __host__
724 double rocrand_uniform_double(rocrand_state_threefry2x32_20* state)
725{
726 return rocrand_device::detail::uniform_distribution_double(rocrand(state));
727}
728
741__forceinline__ __device__ __host__
742 float rocrand_uniform(rocrand_state_threefry2x64_20* state)
743{
744 return rocrand_device::detail::uniform_distribution(rocrand(state));
745}
746
762__forceinline__ __device__ __host__
763 double rocrand_uniform_double(rocrand_state_threefry2x64_20* state)
764{
765 return rocrand_device::detail::uniform_distribution_double(rocrand(state));
766}
767
780__forceinline__ __device__ __host__
781 float rocrand_uniform(rocrand_state_threefry4x32_20* state)
782{
783 return rocrand_device::detail::uniform_distribution(rocrand(state));
784}
785
801__forceinline__ __device__ __host__
802 double rocrand_uniform_double(rocrand_state_threefry4x32_20* state)
803{
804 return rocrand_device::detail::uniform_distribution_double(rocrand(state));
805}
806
819__forceinline__ __device__ __host__
820 float rocrand_uniform(rocrand_state_threefry4x64_20* state)
821{
822 return rocrand_device::detail::uniform_distribution(rocrand(state));
823}
824
840__forceinline__ __device__ __host__
841 double rocrand_uniform_double(rocrand_state_threefry4x64_20* state)
842{
843 return rocrand_device::detail::uniform_distribution_double(rocrand(state));
844}
845 // end of group rocranddevice
847
848#endif // ROCRAND_UNIFORM_H_
rocrand_uniform4
__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:281
rocrand_uniform_double4
__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:338
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:379
rocrand_uniform_double
__forceinline__ __device__ __host__ double rocrand_uniform_double(rocrand_state_philox4x32_10 *state)
Returns a uniformly distributed random double value from (0; 1] range.
Definition rocrand_uniform.h:299
rocrand_uniform
__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:239
rocrand_uniform2
__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:257
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_uniform_double2
__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:320