Ctrl+K

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

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

API library: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-rocrand/checkouts/latest/library/include/rocrand/rocrand_uniform.h Source File
API library
rocrand_uniform.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 
26 #ifndef ROCRAND_UNIFORM_H_
27 #define ROCRAND_UNIFORM_H_
28 
30 #ifndef FQUALIFIERS
31 #define FQUALIFIERS __forceinline__ __device__
32 #endif // FQUALIFIERS
33 
34 #include "rocrand/rocrand_lfsr113.h"
35 #include "rocrand/rocrand_mrg31k3p.h"
36 #include "rocrand/rocrand_mrg32k3a.h"
37 #include "rocrand/rocrand_mtgp32.h"
38 #include "rocrand/rocrand_philox4x32_10.h"
39 #include "rocrand/rocrand_scrambled_sobol32.h"
40 #include "rocrand/rocrand_scrambled_sobol64.h"
41 #include "rocrand/rocrand_sobol32.h"
42 #include "rocrand/rocrand_sobol64.h"
43 #include "rocrand/rocrand_threefry2x32_20.h"
44 #include "rocrand/rocrand_threefry2x64_20.h"
45 #include "rocrand/rocrand_threefry4x32_20.h"
46 #include "rocrand/rocrand_threefry4x64_20.h"
47 #include "rocrand/rocrand_xorwow.h"
48 
49 #include "rocrand/rocrand_common.h"
50 
51 namespace rocrand_device {
52 namespace detail {
53 
54 struct two_uints
55 {
56  unsigned int x;
57  unsigned int y;
58 };
59 
60 union two_uints_to_ulong
61 {
62  two_uints uint2_value;
63  unsigned long long int ulong_value;
64 };
65 
66 // For unsigned integer between 0 and UINT_MAX, returns value between
67 // 0.0f and 1.0f, excluding 0.0f and including 1.0f.
68 FQUALIFIERS
69 float uniform_distribution(unsigned int v)
70 {
71  return ROCRAND_2POW32_INV + (v * ROCRAND_2POW32_INV);
72 }
73 
74 // For unsigned integer between 0 and ULLONG_MAX, returns value between
75 // 0.0f and 1.0f, excluding 0.0f and including 1.0f.
76 FQUALIFIERS
77 float uniform_distribution(unsigned long long int v)
78 {
79  return ROCRAND_2POW32_INV + (v >> 32) * ROCRAND_2POW32_INV;
80 }
81 
82 FQUALIFIERS
83 float4 uniform_distribution4(uint4 v)
84 {
85  return float4 {
86  ROCRAND_2POW32_INV + (v.x * ROCRAND_2POW32_INV),
87  ROCRAND_2POW32_INV + (v.y * ROCRAND_2POW32_INV),
88  ROCRAND_2POW32_INV + (v.z * ROCRAND_2POW32_INV),
89  ROCRAND_2POW32_INV + (v.w * ROCRAND_2POW32_INV)
90  };
91 }
92 
93 FQUALIFIERS float4 uniform_distribution4(ulonglong4 v)
94 {
95  return float4{ROCRAND_2POW64_INV + (v.x * ROCRAND_2POW64_INV),
96  ROCRAND_2POW64_INV + (v.y * ROCRAND_2POW64_INV),
97  ROCRAND_2POW64_INV + (v.z * ROCRAND_2POW64_INV),
98  ROCRAND_2POW64_INV + (v.w * ROCRAND_2POW64_INV)};
99 }
100 
101 // For unsigned integer between 0 and UINT_MAX, returns value between
102 // 0.0 and 1.0, excluding 0.0 and including 1.0.
103 FQUALIFIERS
104 double uniform_distribution_double(unsigned int v)
105 {
106  return ROCRAND_2POW32_INV_DOUBLE + (v * ROCRAND_2POW32_INV_DOUBLE);
107 }
108 
109 FQUALIFIERS
110 double uniform_distribution_double(unsigned int v1, unsigned int v2)
111 {
112  two_uints_to_ulong v;
113  v.uint2_value.x = v1;
114  v.uint2_value.y = (v2 >> 11);
115  return ROCRAND_2POW53_INV_DOUBLE + (v.ulong_value * ROCRAND_2POW53_INV_DOUBLE);
116 }
117 
118 FQUALIFIERS
119 double uniform_distribution_double(unsigned long long int v)
120 {
121  return ROCRAND_2POW53_INV_DOUBLE + (
122  // 2^53 is the biggest int that can be stored in double, such
123  // that it and all smaller integers can be stored in double
124  (v >> 11) * ROCRAND_2POW53_INV_DOUBLE
125  );
126 }
127 
128 FQUALIFIERS
129 double2 uniform_distribution_double2(uint4 v)
130 {
131  return double2 {
132  uniform_distribution_double(v.x, v.y),
133  uniform_distribution_double(v.z, v.w)
134  };
135 }
136 
137 FQUALIFIERS
138 double4 uniform_distribution_double4(uint4 v1, uint4 v2)
139 {
140  return double4 {
141  uniform_distribution_double(v1.x, v1.y),
142  uniform_distribution_double(v1.z, v1.w),
143  uniform_distribution_double(v2.x, v2.y),
144  uniform_distribution_double(v2.z, v2.w)
145  };
146 }
147 
148 FQUALIFIERS double2 uniform_distribution_double2(ulonglong2 v)
149 {
150  return double2{uniform_distribution_double(v.x), uniform_distribution_double(v.y)};
151 }
152 
153 FQUALIFIERS double2 uniform_distribution_double2(ulonglong4 v)
154 {
155  return double2{uniform_distribution_double(v.x), uniform_distribution_double(v.y)};
156 }
157 
158 FQUALIFIERS double4 uniform_distribution_double4(ulonglong4 v)
159 {
160  return double4{uniform_distribution_double(v.x),
161  uniform_distribution_double(v.z),
162  uniform_distribution_double(v.x),
163  uniform_distribution_double(v.z)};
164 }
165 
166 FQUALIFIERS
167 __half uniform_distribution_half(unsigned short v)
168 {
169  return __float2half(ROCRAND_2POW16_INV + (v * ROCRAND_2POW16_INV));
170 }
171 
172 // For an unsigned integer produced by an MRG-based engine, returns a value
173 // in range [0, UINT32_MAX].
174 template<typename state_type>
175 FQUALIFIERS unsigned int mrg_uniform_distribution_uint(unsigned int v) = delete;
176 
177 template<>
178 FQUALIFIERS unsigned int mrg_uniform_distribution_uint<rocrand_state_mrg31k3p>(unsigned int v)
179 {
180  return static_cast<unsigned int>((v - 1) * ROCRAND_MRG31K3P_UINT32_NORM);
181 }
182 
183 template<>
184 FQUALIFIERS unsigned int mrg_uniform_distribution_uint<rocrand_state_mrg32k3a>(unsigned int v)
185 {
186  return static_cast<unsigned int>((v - 1) * ROCRAND_MRG32K3A_UINT_NORM);
187 }
188 
189 // For an unsigned integer produced by an MRG-based engine, returns value between
190 // 0.0f and 1.0f, excluding 0.0f and including 1.0f.
191 template<typename state_type>
192 FQUALIFIERS float mrg_uniform_distribution(unsigned int v) = delete;
193 
194 template<>
195 FQUALIFIERS float mrg_uniform_distribution<rocrand_state_mrg31k3p>(unsigned int v)
196 {
197  double ret = static_cast<double>(v) * ROCRAND_MRG31K3P_NORM_DOUBLE;
198  return static_cast<float>(ret);
199 }
200 
201 template<>
202 FQUALIFIERS float mrg_uniform_distribution<rocrand_state_mrg32k3a>(unsigned int v)
203 {
204  double ret = static_cast<double>(v) * ROCRAND_MRG32K3A_NORM_DOUBLE;
205  return static_cast<float>(ret);
206 }
207 
208 // For an unsigned integer produced by an MRG generator, returns value between
209 // 0.0 and 1.0, excluding 0.0 and including 1.0.
210 template<typename state_type>
211 FQUALIFIERS double mrg_uniform_distribution_double(unsigned int v) = delete;
212 
213 template<>
214 FQUALIFIERS double mrg_uniform_distribution_double<rocrand_state_mrg31k3p>(unsigned int v)
215 {
216  double ret = static_cast<double>(v) * ROCRAND_MRG31K3P_NORM_DOUBLE;
217  return ret;
218 }
219 
220 template<>
221 FQUALIFIERS double mrg_uniform_distribution_double<rocrand_state_mrg32k3a>(unsigned int v)
222 {
223  double ret = static_cast<double>(v) * ROCRAND_MRG32K3A_NORM_DOUBLE;
224  return ret;
225 }
226 
227 } // end namespace detail
228 } // end namespace rocrand_device
229 
242 FQUALIFIERS
243 float rocrand_uniform(rocrand_state_philox4x32_10 * state)
244 {
245  return rocrand_device::detail::uniform_distribution(rocrand(state));
246 }
247 
260 FQUALIFIERS
261 float2 rocrand_uniform2(rocrand_state_philox4x32_10 * state)
262 {
263  auto state1 = rocrand(state);
264  auto state2 = rocrand(state);
265 
266  return float2 {
267  rocrand_device::detail::uniform_distribution(state1),
268  rocrand_device::detail::uniform_distribution(state2)
269  };
270 }
271 
284 FQUALIFIERS
285 float4 rocrand_uniform4(rocrand_state_philox4x32_10 * state)
286 {
287  return rocrand_device::detail::uniform_distribution4(rocrand4(state));
288 }
289 
302 FQUALIFIERS
303 double rocrand_uniform_double(rocrand_state_philox4x32_10 * state)
304 {
305  auto state1 = rocrand(state);
306  auto state2 = rocrand(state);
307 
308  return rocrand_device::detail::uniform_distribution_double(state1, state2);
309 }
310 
323 FQUALIFIERS
324 double2 rocrand_uniform_double2(rocrand_state_philox4x32_10 * state)
325 {
326  return rocrand_device::detail::uniform_distribution_double2(rocrand4(state));
327 }
328 
341 FQUALIFIERS
342 double4 rocrand_uniform_double4(rocrand_state_philox4x32_10 * state)
343 {
344  return rocrand_device::detail::uniform_distribution_double4(rocrand4(state), rocrand4(state));
345 }
346 
359 FQUALIFIERS float rocrand_uniform(rocrand_state_mrg31k3p* state)
360 {
361  return rocrand_device::detail::mrg_uniform_distribution<rocrand_state_mrg31k3p>(state->next());
362 }
363 
379 FQUALIFIERS double rocrand_uniform_double(rocrand_state_mrg31k3p* state)
380 {
381  return rocrand_device::detail::mrg_uniform_distribution_double<rocrand_state_mrg31k3p>(
382  state->next());
383 }
384 
397 FQUALIFIERS
398 float rocrand_uniform(rocrand_state_mrg32k3a * state)
399 {
400  return rocrand_device::detail::mrg_uniform_distribution<rocrand_state_mrg32k3a>(state->next());
401 }
402 
418 FQUALIFIERS
419 double rocrand_uniform_double(rocrand_state_mrg32k3a * state)
420 {
421  return rocrand_device::detail::mrg_uniform_distribution_double<rocrand_state_mrg32k3a>(
422  state->next());
423 }
424 
437 FQUALIFIERS
438 float rocrand_uniform(rocrand_state_xorwow * state)
439 {
440  return rocrand_device::detail::uniform_distribution(rocrand(state));
441 }
442 
455 FQUALIFIERS
456 double rocrand_uniform_double(rocrand_state_xorwow * state)
457 {
458  auto state1 = rocrand(state);
459  auto state2 = rocrand(state);
460 
461  return rocrand_device::detail::uniform_distribution_double(state1, state2);
462 }
463 
476 FQUALIFIERS
477 float rocrand_uniform(rocrand_state_mtgp32 * state)
478 {
479  return rocrand_device::detail::uniform_distribution(rocrand(state));
480 }
481 
497 FQUALIFIERS
498 double rocrand_uniform_double(rocrand_state_mtgp32 * state)
499 {
500  return rocrand_device::detail::uniform_distribution_double(rocrand(state));
501 }
502 
515 FQUALIFIERS
516 float rocrand_uniform(rocrand_state_sobol32 * state)
517 {
518  return rocrand_device::detail::uniform_distribution(rocrand(state));
519 }
520 
536 FQUALIFIERS
537 double rocrand_uniform_double(rocrand_state_sobol32 * state)
538 {
539  return rocrand_device::detail::uniform_distribution_double(rocrand(state));
540 }
541 
554 FQUALIFIERS
555 float rocrand_uniform(rocrand_state_scrambled_sobol32* state)
556 {
557  return rocrand_device::detail::uniform_distribution(rocrand(state));
558 }
559 
575 FQUALIFIERS
576 double rocrand_uniform_double(rocrand_state_scrambled_sobol32* state)
577 {
578  return rocrand_device::detail::uniform_distribution_double(rocrand(state));
579 }
580 
593 FQUALIFIERS
594 float rocrand_uniform(rocrand_state_sobol64* state)
595 {
596  return rocrand_device::detail::uniform_distribution(rocrand(state));
597 }
598 
611 FQUALIFIERS
612 double rocrand_uniform_double(rocrand_state_sobol64 * state)
613 {
614  return rocrand_device::detail::uniform_distribution_double(rocrand(state));
615 }
616 
629 FQUALIFIERS
630 float rocrand_uniform(rocrand_state_scrambled_sobol64* state)
631 {
632  return rocrand_device::detail::uniform_distribution(rocrand(state));
633 }
634 
647 FQUALIFIERS
648 double rocrand_uniform_double(rocrand_state_scrambled_sobol64* state)
649 {
650  return rocrand_device::detail::uniform_distribution_double(rocrand(state));
651 }
652 
665 FQUALIFIERS
666 float rocrand_uniform(rocrand_state_lfsr113* state)
667 {
668  return rocrand_device::detail::uniform_distribution(rocrand(state));
669 }
670 
686 FQUALIFIERS
687 double rocrand_uniform_double(rocrand_state_lfsr113* state)
688 {
689  return rocrand_device::detail::uniform_distribution_double(rocrand(state));
690 }
691 
704 FQUALIFIERS float rocrand_uniform(rocrand_state_threefry2x32_20* state)
705 {
706  return rocrand_device::detail::uniform_distribution(rocrand(state));
707 }
708 
724 FQUALIFIERS double rocrand_uniform_double(rocrand_state_threefry2x32_20* state)
725 {
726  return rocrand_device::detail::uniform_distribution_double(rocrand(state));
727 }
728 
741 FQUALIFIERS float rocrand_uniform(rocrand_state_threefry2x64_20* state)
742 {
743  return rocrand_device::detail::uniform_distribution(rocrand(state));
744 }
745 
761 FQUALIFIERS double rocrand_uniform_double(rocrand_state_threefry2x64_20* state)
762 {
763  return rocrand_device::detail::uniform_distribution_double(rocrand(state));
764 }
765 
778 FQUALIFIERS float rocrand_uniform(rocrand_state_threefry4x32_20* state)
779 {
780  return rocrand_device::detail::uniform_distribution(rocrand(state));
781 }
782 
798 FQUALIFIERS double rocrand_uniform_double(rocrand_state_threefry4x32_20* state)
799 {
800  return rocrand_device::detail::uniform_distribution_double(rocrand(state));
801 }
802 
815 FQUALIFIERS float rocrand_uniform(rocrand_state_threefry4x64_20* state)
816 {
817  return rocrand_device::detail::uniform_distribution(rocrand(state));
818 }
819 
835 FQUALIFIERS double rocrand_uniform_double(rocrand_state_threefry4x64_20* state)
836 {
837  return rocrand_device::detail::uniform_distribution_double(rocrand(state));
838 }
839  // end of group rocranddevice
841 
842 #endif // ROCRAND_UNIFORM_H_
rocrand_uniform2
FQUALIFIERS float2 rocrand_uniform2(rocrand_state_philox4x32_10 *state)
Returns two uniformly distributed random float values from (0; 1] range.
Definition: rocrand_uniform.h:261
rocrand_uniform_double4
FQUALIFIERS double4 rocrand_uniform_double4(rocrand_state_philox4x32_10 *state)
Returns four uniformly distributed random double values from (0; 1] range.
Definition: rocrand_uniform.h:342
rocrand_uniform_double
FQUALIFIERS double rocrand_uniform_double(rocrand_state_philox4x32_10 *state)
Returns a uniformly distributed random double value from (0; 1] range.
Definition: rocrand_uniform.h:303
rocrand
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
rocrand_uniform
FQUALIFIERS float rocrand_uniform(rocrand_state_philox4x32_10 *state)
Returns a uniformly distributed random float value from (0; 1] range.
Definition: rocrand_uniform.h:243
rocrand4
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
rocrand_uniform4
FQUALIFIERS float4 rocrand_uniform4(rocrand_state_philox4x32_10 *state)
Returns four uniformly distributed random float values from (0; 1] range.
Definition: rocrand_uniform.h:285
rocrand_uniform_double2
FQUALIFIERS double2 rocrand_uniform_double2(rocrand_state_philox4x32_10 *state)
Returns two uniformly distributed random double values from (0; 1] range.
Definition: rocrand_uniform.h:324
FQUALIFIERS
#define FQUALIFIERS
Shorthand for commonly used function qualifiers.
Definition: rocrand_uniform.h:31
Generated by doxygen 1.9.1